| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

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

Classes | Public Types | Public Member Functions | Static Public Attributes | Protected Attributes | List of all members
cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable > Class Template Reference

`#include <mma_base.h>`

## Classes
class | SharedStorage
| Shared storage object needed by threadblock-scoped GEMM. More...

## Public Types
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 | MmaBase (SharedStorage &shared_storage, int thread_idx, int warp_idx, int lane_idx)
| Construct from tensor references. More...

## Static Public Attributes
static int const | kWarpGemmIterations
| Number of warp-level GEMM oeprations. More...
static int const | kStages = Stages
| Number of stages. More...

## Protected Attributes
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...

## Detailed Description

### template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>
class cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >

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

## Member Typedef Documentation

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::Operator = typename Policy::Operator

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::Policy = Policy_

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::Shape = Shape_

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::TensorRefA = TensorRef<typename Operator::ElementA, typename Operator::LayoutA>

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::TensorRefB = TensorRef<typename Operator::ElementB, typename Operator::LayoutB>

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::WarpCount = GemmShape<Shape::kM / WarpGemm::kM, Shape::kN / WarpGemm::kN, Shape::kK / WarpGemm::kK>

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::WarpGemm = typename Policy::Operator::Shape

Shape describing the overall GEMM computed from shared memory by each warp.

## Constructor & Destructor Documentation

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

CUTLASS_DEVICE cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::MmaBase | ( | 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 ID within the threadblock
  thread_idx | ID of warp
  warp_idx | ID of each thread within a warp

## Member Data Documentation

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

int const cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::kStages = Stages
| static

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

int const cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::kWarpGemmIterations
| static

Initial value:
=
(WarpGemm::kK / Operator::Policy::MmaShape::kK)
template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

Operator::IteratorA cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::warp_tile_iterator_A_
| protected

template<typename Shape_, typename Policy_, int Stages, typename Enable = bool>

Operator::IteratorB cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::warp_tile_iterator_B_
| protected

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

- mma_base.h