CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members
cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable > Class Template Reference

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

#include <mma_tensor_op_sm70.h>

Public Types

using Shape = Shape_
 Shape of warp-level matrix operation (concept: GemmShape) More...
 
using ElementA = ElementA_
 Data type of multiplicand A. More...
 
using LayoutA = LayoutA_
 Layout of multiplicand A. More...
 
using ElementB = ElementB_
 Data type of multiplicand B. More...
 
using LayoutB = LayoutB_
 Layout of multiplicand B. More...
 
using ElementC = ElementC_
 Data type of accumulator matrix C. More...
 
using LayoutC = LayoutC_
 Layout of accumulator matrix C. More...
 
using Policy = Policy_
 Shape of the warp in units of thread (concept: MmaLanePolicySimt) More...
 
using OperatorClass = arch::OpClassTensorOp
 Indicates class of matrix operator. More...
 
using InterleavedTileShape = GemmShape< 32, 32, 4 >
 interleaved 32x32 tiles More...
 
using IteratorA = MmaVoltaTensorOpMultiplicandTileIterator< MatrixShape< Shape::kM, Shape::kK >, Operand::kA, ElementA, LayoutA, MatrixShape< Policy::Operator::Shape::kM, Policy::Operator::Shape::kK >, Policy::OpDelta::kRow, kThreadCount >
 Iterates over the A operand in memory. More...
 
using FragmentA = typename IteratorA::Fragment
 Storage for A tile. More...
 
using IteratorB = MmaVoltaTensorOpMultiplicandTileIterator< MatrixShape< Shape::kK, Shape::kN >, Operand::kB, ElementB, LayoutB, MatrixShape< Policy::Operator::Shape::kK, Policy::Operator::Shape::kN >, Policy::OpDelta::kRow, kThreadCount >
 Iterates over the B operand in memory. More...
 
using FragmentB = typename IteratorB::Fragment
 Storage for B tile. More...
 
using IteratorC = MmaVoltaTensorOpAccumulatorTileIterator< MatrixShape< Shape::kM, Shape::kN >, ElementC, LayoutC, typename Policy::Operator::Shape, typename Policy::OpDelta >
 Iterates over the C operand in memory. More...
 
using FragmentC = typename IteratorC::Fragment
 Storage for C tile. More...
 

Public Member Functions

CUTLASS_DEVICE MmaVoltaTensorOp ()
 Ctor. More...
 
CUTLASS_DEVICE void operator() (FragmentC &D, FragmentA const &A, FragmentB const &B, FragmentC const &C, int const &partitionN_idx=0)
 Performs a warp-level matrix multiply-accumulate operation. More...
 

Public Attributes

Policy::Operator mma
 Underlying matrix multiply operator (concept: arch::Mma) More...
 

Static Public Attributes

static int const kThreadCount = 32
 Number of threads participating in warp-level matrix product. More...
 

Member Typedef Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::ElementA = ElementA_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::ElementB = ElementB_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::ElementC = ElementC_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::FragmentA = typename IteratorA::Fragment
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::FragmentB = typename IteratorB::Fragment
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::FragmentC = typename IteratorC::Fragment
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::InterleavedTileShape = GemmShape<32, 32, 4>
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::IteratorA = MmaVoltaTensorOpMultiplicandTileIterator< MatrixShape<Shape::kM, Shape::kK>, Operand::kA, ElementA, LayoutA, MatrixShape< Policy::Operator::Shape::kM, Policy::Operator::Shape::kK >, Policy::OpDelta::kRow, kThreadCount >
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::IteratorB = MmaVoltaTensorOpMultiplicandTileIterator< MatrixShape<Shape::kK, Shape::kN>, Operand::kB, ElementB, LayoutB, MatrixShape< Policy::Operator::Shape::kK, Policy::Operator::Shape::kN >, Policy::OpDelta::kRow, kThreadCount >
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::IteratorC = MmaVoltaTensorOpAccumulatorTileIterator< MatrixShape<Shape::kM, Shape::kN>, ElementC, LayoutC, typename Policy::Operator::Shape, typename Policy::OpDelta >
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::LayoutA = LayoutA_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::LayoutB = LayoutB_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::LayoutC = LayoutC_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::OperatorClass = arch::OpClassTensorOp
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::Policy = Policy_
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
using cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::Shape = Shape_

Constructor & Destructor Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
CUTLASS_DEVICE cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::MmaVoltaTensorOp ( )
inline

Member Function Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::operator() ( FragmentC &  D,
FragmentA const &  A,
FragmentB const &  B,
FragmentC const &  C,
int const &  partitionN_idx = 0 
)
inline

Member Data Documentation

template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
int const cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::kThreadCount = 32
static
template<typename Shape_ , typename ElementA_ , typename LayoutA_ , typename ElementB_ , typename LayoutB_ , typename ElementC_ , typename LayoutC_ , typename Policy_ , typename Enable = bool>
Policy::Operator cutlass::gemm::warp::MmaVoltaTensorOp< Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, Policy_, Enable >::mma
The documentation for this class was generated from the following file:
Generated by   doxygen 1.8.11