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

#include <mma_pipelined.h>

Inheritance diagram for cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >:

[legend]

Collaboration diagram for cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >:

[legend]

Public Types
using	Base = MmaBase< Shape_, Policy_, 2 >
	< Base class More...

using	Shape = Shape_
	Size of the Gemm problem - concept: gemm::GemmShape<> More...

using	IteratorA = IteratorA_
	Iterates over tiles of A operand in global memory. More...

using	IteratorB = IteratorB_
	Iterates over tiles of B operand in global memory. More...

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

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

using	Policy = Policy_
	Policy describing tuning details. More...

using	SmemIteratorA = SmemIteratorA_

using	SmemIteratorB = SmemIteratorB_

using	TransformA = TransformA_

using	TransformB = TransformB_

using	FragmentA = typename IteratorA::Fragment
	Fragment of operand A loaded from global memory. More...

using	FragmentB = typename IteratorB::Fragment
	Fragment of operand B loaded from global memory. More...

using	FragmentC = typename Policy::Operator::FragmentC
	Fragment of accumulator tile. More...

using	Operator = typename Policy::Operator
	Warp-level Mma. More...

Public Types inherited from cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, 2 >
using	Shape = Shape_
	Policy describing tuning details. More...

using	Policy = Policy_

using	Operator = typename Policy::Operator
	Warp-level Mma. More...

using	WarpGemm = typename Policy::Operator::Shape

using	WarpCount = GemmShape< Shape::kM/WarpGemm::kM, Shape::kN/WarpGemm::kN, Shape::kK/WarpGemm::kK >
	Shape describing the number of warps filling the CTA. More...

using	TensorRefA = TensorRef< typename Operator::ElementA, typename Operator::LayoutA >
	Tensor reference to the A operand. More...

using	TensorRefB = TensorRef< typename Operator::ElementB, typename Operator::LayoutB >
	Tensor reference to the B operand. More...

Public Member Functions
CUTLASS_DEVICE	MmaPipelined (typename Base::SharedStorage &shared_storage, int thread_idx, int warp_idx, int lane_idx)
	Construct from tensor references. More...

CUTLASS_DEVICE void	operator() (int gemm_k_iterations, FragmentC &accum, IteratorA iterator_A, IteratorB iterator_B, FragmentC const &src_accum, TransformA transform_A=TransformA(), TransformB transform_B=TransformB())
	Perform a threadblock-scoped matrix multiply-accumulate. More...

Public Member Functions inherited from cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, 2 >
CUTLASS_DEVICE	MmaBase (SharedStorage &shared_storage, int thread_idx, int warp_idx, int lane_idx)
	Construct from tensor references. More...

Protected Attributes
SmemIteratorA	smem_iterator_A_
	Iterator to write threadblock-scoped tile of A operand to shared memory. More...

SmemIteratorB	smem_iterator_B_
	Iterator to write threadblock-scoped tile of B operand to shared memory. More...

Protected Attributes inherited from cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, 2 >
Operator::IteratorA	warp_tile_iterator_A_
	Iterator to load a warp-scoped tile of A operand from shared memory. More...

Operator::IteratorB	warp_tile_iterator_B_
	Iterator to load a warp-scoped tile of B operand from shared memory. More...

Additional Inherited Members
Static Public Attributes inherited from cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, 2 >
static int const	kWarpGemmIterations
	Number of warp-level GEMM oeprations. More...

static int const	kStages
	Number of stages. More...

Member Typedef Documentation

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::Base = MmaBase<Shape_, Policy_, 2>

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::ElementC = ElementC_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::FragmentA = typename IteratorA::Fragment

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::FragmentB = typename IteratorB::Fragment

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::FragmentC = typename Policy::Operator::FragmentC

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::IteratorA = IteratorA_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::IteratorB = IteratorB_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::LayoutC = LayoutC_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::Operator = typename Policy::Operator

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::Policy = Policy_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::Shape = Shape_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::SmemIteratorA = SmemIteratorA_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::SmemIteratorB = SmemIteratorB_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::TransformA = TransformA_

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

using cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::TransformB = TransformB_

Constructor & Destructor Documentation

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

CUTLASS_DEVICE cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::MmaPipelined	(	typename Base::SharedStorage &	shared_storage,
		int	thread_idx,
		int	warp_idx,
		int	lane_idx
	)

inline

Parameters

shared_storage	Shared storage needed for internal use by threadblock-scoped GEMM
thread_idx	ID within the threadblock
warp_idx	ID of warp
lane_idx	ID of each thread within a warp

Member Function Documentation

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

CUTLASS_DEVICE void cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::operator()	(	int	gemm_k_iterations,
		FragmentC &	accum,
		IteratorA	iterator_A,
		IteratorB	iterator_B,
		FragmentC const &	src_accum,
		TransformA	transform_A = `TransformA()`,
		TransformB	transform_B = `TransformB()`
	)

inline

< transformation applied to B fragment

Parameters

gemm_k_iterations	number of iterations of the mainloop
accum	destination accumulator tile
iterator_A	iterator over A operand in global memory
iterator_B	iterator over B operand in global memory
src_accum	source accumulator tile
transform_A	transformation applied to A fragment

Member Data Documentation

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

SmemIteratorA cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::smem_iterator_A_

protected

template<typename Shape_ , typename IteratorA_ , typename SmemIteratorA_ , typename IteratorB_ , typename SmemIteratorB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename TransformA_ = NumericArrayConverter< typename SmemIteratorA_::Element, typename IteratorA_::Element, IteratorA_::Fragment::kElements>, typename TransformB_ = NumericArrayConverter< typename SmemIteratorB_::Element, typename IteratorB_::Element, IteratorB_::Fragment::kElements>, typename Enable = bool>

SmemIteratorB cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >::smem_iterator_B_

protected

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

mma_pipelined.h

Public Types

Public Member Functions

Protected Attributes

Additional Inherited Members

Member Typedef Documentation

Constructor & Destructor Documentation

Member Function Documentation

Member Data Documentation