CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Public Member Functions | Protected Attributes | List of all members
cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable > Class Template Reference

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 >:
Inheritance graph
[legend]
Collaboration diagram for cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable >:
Collaboration graph
[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_storageShared storage needed for internal use by threadblock-scoped GEMM
thread_idxID within the threadblock
warp_idxID of warp
lane_idxID 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_iterationsnumber of iterations of the mainloop
accumdestination accumulator tile
iterator_Aiterator over A operand in global memory
iterator_Biterator over B operand in global memory
src_accumsource accumulator tile
transform_Atransformation 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: