Partial specialization for Turing Integer Matrix Multiply Interleaved layout.

#include <default_gemm.h>

Public Types
using	LayoutA = layout::ColumnMajorInterleaved< InterleavedK >

using	LayoutB = layout::RowMajorInterleaved< InterleavedK >

using	LayoutC = layout::ColumnMajorInterleaved< InterleavedK >

using	ElementAccumulator = int32_t

using	Mma = typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, LayoutC, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, true >::ThreadblockMma
	Define the threadblock-scoped matrix multiply-accumulate. More...

using	Epilogue = typename cutlass::epilogue::threadblock::DefaultInterleavedEpilogueTensorOp< ThreadblockShape, typename Mma::Operator, kPartitionsK, EpilogueOutputOp, 64/sizeof_bits< ElementC >::value, InterleavedK, IsBetaZero >::Epilogue
	Define the epilogue. More...

using	GemmKernel = kernel::Gemm< Mma, Epilogue, ThreadblockSwizzle, SplitKSerial >
	Define the kernel-level GEMM operator. More...

Static Public Attributes
static const int	kPartitionsK = ThreadblockShape::kK / WarpShape::kK

Member Typedef Documentation

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::ElementAccumulator = int32_t

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::Epilogue = typename cutlass::epilogue::threadblock:: DefaultInterleavedEpilogueTensorOp< ThreadblockShape, typename Mma::Operator, kPartitionsK, EpilogueOutputOp, 64 / sizeof_bits<ElementC>::value, InterleavedK, IsBetaZero>::Epilogue

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::LayoutA = layout::ColumnMajorInterleaved<InterleavedK>

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::LayoutB = layout::RowMajorInterleaved<InterleavedK>

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::LayoutC = layout::ColumnMajorInterleaved<InterleavedK>

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

using cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::Mma = typename cutlass::gemm::threadblock::DefaultMma< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, LayoutC, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, true>::ThreadblockMma

Member Data Documentation

template<typename ElementA , int kAlignmentA, typename ElementB , int kAlignmentB, typename ElementC , typename ThreadblockShape , typename WarpShape , typename InstructionShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , int InterleavedK, bool SplitKSerial, typename Operator , bool IsBetaZero>

const int cutlass::gemm::kernel::DefaultGemm< ElementA, layout::ColumnMajorInterleaved< InterleavedK >, kAlignmentA, ElementB, layout::RowMajorInterleaved< InterleavedK >, kAlignmentB, ElementC, layout::ColumnMajorInterleaved< InterleavedK >, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero >::kPartitionsK = ThreadblockShape::kK / WarpShape::kK

static

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

default_gemm.h

Public Types

Static Public Attributes

Member Typedef Documentation

Member Data Documentation