ContextQMD
Back to Cutlass

CUTLASS: cutlass::gemm::threadblock::MmaPipelined< Shape_, IteratorA_, SmemIteratorA_, IteratorB_, SmemIteratorB_, ElementC_, LayoutC_, Policy_, TransformA_, TransformB_, Enable > Class Template Reference

docs/classcutlass_1_1gemm_1_1threadblock_1_1MmaPipelined.html

4.4.229.2 KB
Original Source

| | 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](classcutlass_1_1gemm_1_1threadblock_1_1MmaPipelined inherit graph.png) [legend]

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

![Collaboration graph](classcutlass_1_1gemm_1_1threadblock_1_1MmaPipelined coll graph.png) [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:

Generated by 1.8.11