CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Static Public Attributes | List of all members
cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ > Struct Template Reference

#include <default_mma_core_simt.h>

Public Types

using Shape = Shape_
 
using WarpShape = WarpShape_
 
using InstructionShape = GemmShape< 1, 1, 1 >
 
using ElementA = ElementA_
 
using LayoutA = layout::ColumnMajor
 
using ElementB = ElementB_
 
using LayoutB = layout::RowMajor
 
using ElementC = ElementC_
 
using LayoutC = LayoutC_
 
using OperatorClass = arch::OpClassSimt
 
using Operator = Operator_
 Default Operator. More...
 
using WarpCount = GemmShape< Shape::kM/WarpShape::kM, Shape::kN/WarpShape::kN, PartitionsK >
 Number of warps present. More...
 
using SmemLayoutA = layout::ColumnMajor
 
using SmemLayoutB = layout::RowMajor
 
using IteratorThreadMapA = transform::PitchLinearStripminedThreadMap< layout::PitchLinearShape< Shape::kM, Shape::kK >, kThreads, kElementsPerAccess >
 ThreadMap of iterator A. More...
 
using SmemIteratorA = transform::threadblock::RegularTileIterator< MatrixShape< Shape::kM, Shape::kK >, ElementA, SmemLayoutA, 1, IteratorThreadMapA >
 Shared memory iterator to A operand. More...
 
using IteratorThreadMapB = transform::PitchLinearStripminedThreadMap< layout::PitchLinearShape< Shape::kN, Shape::kK >, kThreads, kElementsPerAccess >
 Policy of iterator B. More...
 
using SmemIteratorB = transform::threadblock::RegularTileIterator< MatrixShape< Shape::kK, Shape::kN >, ElementB, SmemLayoutB, 0, IteratorThreadMapB >
 Shared memory iterator to B operand. More...
 
using LaneMmaShape = cutlass::gemm::GemmShape< LaneM, LaneN, 1 >
 
using Policy = cutlass::gemm::warp::MmaSimtPolicy< cutlass::MatrixShape< WarpNumThreadsM, WarpNumThreadsN >, cutlass::layout::RowMajorInterleaved< LaneLayout >, LaneMmaShape >
 
using MmaWarpSimt = cutlass::gemm::warp::MmaSimt< WarpShape, ElementA, SmemLayoutA, ElementB, SmemLayoutB, ElementC, LayoutC, Policy >
 
using MmaPolicy = MmaPolicy< MmaWarpSimt, MatrixShape< 0, 0 >, MatrixShape< 0, 0 >, WarpCount::kK >
 Used for partial specialization. More...
 

Static Public Attributes

static int const PartitionsK = Shape::kK / WarpShape::kK
 
static int const kWarpSize = warp::WarpSize<arch::OpClassSimt>::value
 Number of threads per warp. More...
 
static int const kThreads = WarpCount::kCount * kWarpSize
 Number of threads total. More...
 
static int const kElementsPerAccess = 1
 
static const int WarpNumThreadsM = detail::simt_get_warp_threads_m<WarpShape>()
 
static const int WarpNumThreadsN = kWarpSize / WarpNumThreadsM
 
static const int ThreadTileM = WarpShape::kM / WarpNumThreadsM
 
static const int ThreadTileN = WarpShape::kN / WarpNumThreadsN
 
static const int LaneLayout = ThreadTileM > 4 && ThreadTileN > 4 ? 2 : 1
 
static const int numElementsA = 128 / sizeof_bits<ElementA>::value
 
static const int numElementsB = 128 / sizeof_bits<ElementB>::value
 
static const int LaneM = cutlass::const_min(numElementsA, ThreadTileM)
 
static const int LaneN = cutlass::const_min(numElementsB, ThreadTileN)
 

Detailed Description

template<typename Shape_, typename WarpShape_, typename ElementA_, typename ElementB_, typename ElementC_, typename LayoutC_, typename Operator_>
struct cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >

Partial specialization:

A: column-major B: row-major Operator: simt class

This uses the default warp-level operator given tile sizes

Member Typedef Documentation

template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::ElementA = ElementA_
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::ElementB = ElementB_
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::ElementC = ElementC_
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::InstructionShape = GemmShape<1, 1, 1>
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::IteratorThreadMapA = transform::PitchLinearStripminedThreadMap< layout::PitchLinearShape<Shape::kM, Shape::kK>, kThreads, kElementsPerAccess >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::IteratorThreadMapB = transform::PitchLinearStripminedThreadMap< layout::PitchLinearShape<Shape::kN, Shape::kK>, kThreads, kElementsPerAccess >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LaneMmaShape = cutlass::gemm::GemmShape< LaneM, LaneN, 1>
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LayoutA = layout::ColumnMajor
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LayoutB = layout::RowMajor
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LayoutC = LayoutC_
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::MmaPolicy = MmaPolicy< MmaWarpSimt, MatrixShape<0, 0>, MatrixShape<0, 0>, WarpCount::kK >

Policy used to define MmaPipelined

template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::MmaWarpSimt = cutlass::gemm::warp::MmaSimt< WarpShape, ElementA, SmemLayoutA, ElementB, SmemLayoutB, ElementC, LayoutC, Policy >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::Operator = Operator_
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::OperatorClass = arch::OpClassSimt
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::Policy = cutlass::gemm::warp::MmaSimtPolicy< cutlass::MatrixShape<WarpNumThreadsM, WarpNumThreadsN>, cutlass::layout::RowMajorInterleaved<LaneLayout>, LaneMmaShape >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::Shape = Shape_
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::SmemIteratorA = transform::threadblock::RegularTileIterator< MatrixShape<Shape::kM, Shape::kK>, ElementA, SmemLayoutA, 1, IteratorThreadMapA >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::SmemIteratorB = transform::threadblock::RegularTileIterator< MatrixShape<Shape::kK, Shape::kN>, ElementB, SmemLayoutB, 0, IteratorThreadMapB >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::SmemLayoutA = layout::ColumnMajor
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::SmemLayoutB = layout::RowMajor
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::WarpCount = GemmShape< Shape::kM / WarpShape::kM, Shape::kN / WarpShape::kN, PartitionsK >
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
using cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::WarpShape = WarpShape_

Member Data Documentation

template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
int const cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::kElementsPerAccess = 1
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
int const cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::kThreads = WarpCount::kCount * kWarpSize
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
int const cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::kWarpSize = warp::WarpSize<arch::OpClassSimt>::value
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LaneLayout = ThreadTileM > 4 && ThreadTileN > 4 ? 2 : 1
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LaneM = cutlass::const_min(numElementsA, ThreadTileM)
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::LaneN = cutlass::const_min(numElementsB, ThreadTileN)
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::numElementsA = 128 / sizeof_bits<ElementA>::value
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::numElementsB = 128 / sizeof_bits<ElementB>::value
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
int const cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::PartitionsK = Shape::kK / WarpShape::kK
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::ThreadTileM = WarpShape::kM / WarpNumThreadsM
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::ThreadTileN = WarpShape::kN / WarpNumThreadsN
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::WarpNumThreadsM = detail::simt_get_warp_threads_m<WarpShape>()
static
template<typename Shape_ , typename WarpShape_ , typename ElementA_ , typename ElementB_ , typename ElementC_ , typename LayoutC_ , typename Operator_ >
const int cutlass::gemm::threadblock::DefaultMmaCore< Shape_, WarpShape_, GemmShape< 1, 1, 1 >, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ >::WarpNumThreadsN = kWarpSize / WarpNumThreadsM
static

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