| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- Class List
- Class Index
- Class Hierarchy
- Class Members

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 >:
[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