cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage Class Reference

Shared storage object needed by threadblock-scoped GEMM.

#include <mma_base.h>

Collaboration diagram for cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage:

Public Types
using	ShapeA = MatrixShape< Shape::kM+Policy::SmemPaddingA::kRow, Shape::kK *kStages+Policy::SmemPaddingA::kColumn >
	Shape of the A matrix operand in shared memory. More...
using	ShapeB = MatrixShape< Shape::kK *kStages+Policy::SmemPaddingB::kRow, Shape::kN+Policy::SmemPaddingB::kColumn >
	Shape of the B matrix operand in shared memory. More...

Public Member Functions
CUTLASS_HOST_DEVICE TensorRefA	operand_A_ref ()
	Returns a TensorRef to the A operand. More...
CUTLASS_HOST_DEVICE TensorRefB	operand_B_ref ()
	Returns a TensorRef to the B operand. More...

Static Public Member Functions
static CUTLASS_DEVICE Operator::LayoutA	LayoutA ()
	Returns a layout object for the A matrix. More...
static CUTLASS_HOST_DEVICE Operator::LayoutB	LayoutB ()
	Returns a layout object for the B matrix. More...

Public Attributes
AlignedBuffer< typename Operator::ElementA, ShapeA::kCount >	operand_A
	Buffer for A operand. More...
AlignedBuffer< typename Operator::ElementB, ShapeB::kCount >	operand_B
	Buffer for B operand. More...

Member Typedef Documentation

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

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::ShapeA = MatrixShape<Shape::kM + Policy::SmemPaddingA::kRow, Shape::kK * kStages + Policy::SmemPaddingA::kColumn>

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

using cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::ShapeB = MatrixShape<Shape::kK * kStages + Policy::SmemPaddingB::kRow, Shape::kN + Policy::SmemPaddingB::kColumn>

Member Function Documentation

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

static CUTLASS_DEVICE Operator::LayoutA cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::LayoutA ( )

inlinestatic

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

static CUTLASS_HOST_DEVICE Operator::LayoutB cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::LayoutB ( )

inlinestatic

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

CUTLASS_HOST_DEVICE TensorRefA cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::operand_A_ref ( )

inline

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

CUTLASS_HOST_DEVICE TensorRefB cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::operand_B_ref ( )

inline

Member Data Documentation

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

AlignedBuffer<typename Operator::ElementA, ShapeA::kCount> cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::operand_A

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

AlignedBuffer<typename Operator::ElementB, ShapeB::kCount> cutlass::gemm::threadblock::MmaBase< Shape_, Policy_, Stages, Enable >::SharedStorage::operand_B

---

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

• mma_base.h