CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Static Public Attributes | List of all members
cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false > Struct Template Reference

#include <default_mma.h>

Public Types

using InstructionShape = GemmShape< 1, 1, 4 >
 
using ElementA = int8_t
 
using ElementB = int8_t
 
using OperatorClass = arch::OpClassSimt
 
using MmaCore = typename cutlass::gemm::threadblock::DefaultMmaCore< ThreadblockShape, WarpShape, InstructionShape, ElementA, LayoutA, ElementB, LayoutB, ElementAccumulator, layout::RowMajor, OperatorClass, 2, Operator >
 
using IteratorA = cutlass::transform::threadblock::PredicatedTileIterator2dThreadTile< cutlass::MatrixShape< MmaCore::Shape::kM, MmaCore::Shape::kK >, ElementA, LayoutA, 1, typename MmaCore::IteratorThreadMapA, transposeA >
 
using IteratorB = cutlass::transform::threadblock::PredicatedTileIterator2dThreadTile< cutlass::MatrixShape< MmaCore::Shape::kK, MmaCore::Shape::kN >, ElementB, LayoutB, 0, typename MmaCore::IteratorThreadMapB, transposeB >
 
using ThreadblockMma = cutlass::gemm::threadblock::MmaPipelined< typename MmaCore::Shape, IteratorA, typename MmaCore::SmemIteratorA, IteratorB, typename MmaCore::SmemIteratorB, ElementAccumulator, layout::RowMajor, typename MmaCore::MmaPolicy >
 

Static Public Attributes

static const bool transposeA = cutlass::platform::is_same< LayoutA, layout::ColumnMajor >::value
 
static const bool transposeB = cutlass::platform::is_same< LayoutB, layout::RowMajor >::value
 

Detailed Description

template<typename LayoutA, int kAlignmentA, typename LayoutB, int kAlignmentB, typename ElementAccumulator, typename ArchTag, typename ThreadblockShape, typename Operator, typename WarpShape>
struct cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >

Specialization for SIMT IDP4A Kernels

Member Typedef Documentation

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::ElementA = int8_t
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::ElementB = int8_t
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::InstructionShape = GemmShape<1, 1, 4>
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::IteratorA = cutlass::transform::threadblock::PredicatedTileIterator2dThreadTile< cutlass::MatrixShape<MmaCore::Shape::kM, MmaCore::Shape::kK>, ElementA, LayoutA, 1, typename MmaCore::IteratorThreadMapA, transposeA>
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::IteratorB = cutlass::transform::threadblock::PredicatedTileIterator2dThreadTile< cutlass::MatrixShape<MmaCore::Shape::kK, MmaCore::Shape::kN>, ElementB, LayoutB, 0, typename MmaCore::IteratorThreadMapB, transposeB>
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::MmaCore = typename cutlass::gemm::threadblock::DefaultMmaCore< ThreadblockShape, WarpShape, InstructionShape, ElementA, LayoutA, ElementB, LayoutB, ElementAccumulator, layout::RowMajor, OperatorClass, 2, Operator>
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::OperatorClass = arch::OpClassSimt
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
using cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::ThreadblockMma = cutlass::gemm::threadblock::MmaPipelined< typename MmaCore::Shape, IteratorA, typename MmaCore::SmemIteratorA, IteratorB, typename MmaCore::SmemIteratorB, ElementAccumulator, layout::RowMajor, typename MmaCore::MmaPolicy>

Member Data Documentation

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
const bool cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::transposeA = cutlass::platform::is_same< LayoutA, layout::ColumnMajor >::value
static
template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename ElementAccumulator , typename ArchTag , typename ThreadblockShape , typename Operator , typename WarpShape >
const bool cutlass::gemm::threadblock::DefaultMma< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, 2, Operator, false >::transposeB = cutlass::platform::is_same< LayoutB, layout::RowMajor >::value
static

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