CUTLASS: cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable > Class Template Reference

CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers

Structure to compute the matrix product targeting CUDA cores and SIMT math instructions.

#include <mma_simt.h>

Public Types
using	Shape = Shape_
	Shape of warp-level matrix operation (concept: GemmShape) More...

using	ElementA = ElementA_
	Data type of multiplicand A. More...

using	LayoutA = LayoutA_
	Layout of multiplicand A. More...

using	ElementB = ElementB_
	Data type of multiplicand B. More...

using	LayoutB = LayoutB_
	Layout of multiplicand B. More...

using	ElementC = ElementC_
	Data type of accumulator matrix C. More...

using	LayoutC = LayoutC_
	Layout of accumulator matrix C. More...

using	Policy = Policy_
	Shape of the warp in units of thread (concept: MmaLanePolicySimt) More...

using	OperatorClass = arch::OpClassSimt
	Indicates class of matrix operator. More...

using	ThreadLayoutA = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved< 4 >, LayoutA >::value, layout::ColumnMajor, typename platform::conditional< platform::is_same< layout::RowMajorInterleaved< 4 >, LayoutA >::value, layout::RowMajor, LayoutA >::type >::type

using	ThreadLayoutB = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved< 4 >, LayoutB >::value, layout::ColumnMajor, typename platform::conditional< platform::is_same< layout::RowMajorInterleaved< 4 >, LayoutB >::value, layout::RowMajor, LayoutB >::type >::type

using	dp4a_type = typename platform::conditional< use_dp4a, int8_t, bool >::type

using	ThreadMma = thread::Mma< GemmShape< Shape::kM/Policy::WarpShape::kRow, Shape::kN/Policy::WarpShape::kColumn, Policy::LaneMmaShape::kK >, ElementA, ThreadLayoutA, ElementB, ThreadLayoutB, ElementC, LayoutC, arch::OpMultiplyAdd, dp4a_type >
	Thread-level matrix multiply accumulate operator. More...

using	IteratorA = MmaSimtTileIterator< MatrixShape< Shape::kM, Policy::LaneMmaShape::kK >, Operand::kA, ElementA, LayoutA, Policy, PartitionsK, Shape::kK >
	Iterates over the A operand in memory. More...

using	FragmentA = typename IteratorA::Fragment
	Storage for A tile. More...

using	IteratorB = MmaSimtTileIterator< MatrixShape< Policy::LaneMmaShape::kK, Shape::kN >, Operand::kB, ElementB, LayoutB, Policy, PartitionsK, Shape::kK >
	Iterates over the B operand in memory. More...

using	FragmentB = typename IteratorB::Fragment
	Storage for B tile. More...

using	IteratorC = MmaSimtTileIterator< MatrixShape< Shape::kM, Shape::kN >, Operand::kC, ElementC, LayoutC, Policy >
	Iterates over the C operand in memory. More...

using	FragmentC = typename ThreadMma::FragmentC
	Storage for C tile. More...

Public Member Functions
CUTLASS_DEVICE	MmaSimt ()
	Ctor. More...

CUTLASS_DEVICE void	operator() (FragmentC &d, FragmentA const &a, FragmentB const &b, FragmentC const &c, int group_idx=0) const
	Performs a warp-level matrix multiply-accumulate operation. More...

Static Public Attributes
static constexpr bool	use_dp4a

Member Typedef Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::dp4a_type = typename platform::conditional< use_dp4a , int8_t, bool >::type

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ElementA = ElementA_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ElementB = ElementB_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ElementC = ElementC_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::FragmentA = typename IteratorA::Fragment

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::FragmentB = typename IteratorB::Fragment

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::FragmentC = typename ThreadMma::FragmentC

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::IteratorA = MmaSimtTileIterator< MatrixShape<Shape::kM, Policy::LaneMmaShape::kK>, Operand::kA, ElementA, LayoutA, Policy, PartitionsK, Shape::kK >

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::IteratorB = MmaSimtTileIterator< MatrixShape<Policy::LaneMmaShape::kK, Shape::kN>, Operand::kB, ElementB, LayoutB, Policy, PartitionsK, Shape::kK >

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::IteratorC = MmaSimtTileIterator< MatrixShape<Shape::kM, Shape::kN>, Operand::kC, ElementC, LayoutC, Policy >

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::LayoutA = LayoutA_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::LayoutB = LayoutB_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::LayoutC = LayoutC_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::OperatorClass = arch::OpClassSimt

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::Policy = Policy_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::Shape = Shape_

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ThreadLayoutA = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved<4>, LayoutA >::value, layout::ColumnMajor, typename platform::conditional < platform::is_same< layout::RowMajorInterleaved<4>, LayoutA >::value, layout::RowMajor, LayoutA>::type >::type

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ThreadLayoutB = typename platform::conditional< platform::is_same< layout::ColumnMajorInterleaved<4>, LayoutB >::value, layout::ColumnMajor, typename platform::conditional < platform::is_same< layout::RowMajorInterleaved<4>, LayoutB >::value, layout::RowMajor, LayoutB>::type >::type

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

using cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::ThreadMma = thread::Mma< GemmShape< Shape::kM / Policy::WarpShape::kRow, Shape::kN / Policy::WarpShape::kColumn, Policy::LaneMmaShape::kK>, ElementA, ThreadLayoutA, ElementB, ThreadLayoutB, ElementC, LayoutC, arch::OpMultiplyAdd, dp4a_type >

Constructor & Destructor Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

CUTLASS_DEVICE cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::MmaSimt ( )

inline

Member Function Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

CUTLASS_DEVICE void cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::operator()	(	FragmentC &	d,
		FragmentA const &	a,
		FragmentB const &	b,
		FragmentC const &	c,
		int	group_idx = `0`
	)		const

inline

Member Data Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , int PartitionsK = 1, typename Enable = bool>

constexpr bool cutlass::gemm::warp::MmaSimt< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, PartitionsK, Enable >::use_dp4a

static

Initial value:

= (platform::is_same< layout::ColumnMajorInterleaved<4>, LayoutA>::value || 
                                    platform::is_same< layout::RowMajorInterleaved<4>, LayoutA >::value) && 
                                    platform::is_same< ElementA, int8_t >::value && 
                                    platform::is_same< ElementB, int8_t >::value

The documentation for this class was generated from the following file:

mma_simt.h