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