Partial specialization for SIMT DP4A.

#include <default_gemm.h>

Public Types
using	InstructionShape = GemmShape< 1, 1, 4 >

using	ElementA = int8_t

using	ElementB = int8_t

using	OperatorClass = arch::OpClassSimt

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

using	Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueSimt< ThreadblockShape, typename Mma::Operator, EpilogueOutputOp, kEpilogueElementsPerAccess >::Epilogue
	Define the epilogue. More...

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

Static Public Attributes
static int const	kEpilogueElementsPerAccess = EpilogueOutputOp::kCount

Member Typedef Documentation

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::ElementA = int8_t

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::ElementB = int8_t

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::Epilogue = typename cutlass::epilogue::threadblock::DefaultEpilogueSimt< ThreadblockShape, typename Mma::Operator, EpilogueOutputOp, kEpilogueElementsPerAccess >::Epilogue

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::GemmKernel = kernel::Gemm<Mma, Epilogue, ThreadblockSwizzle, SplitKSerial>

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::InstructionShape = GemmShape<1, 1, 4>

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::Mma = typename cutlass::gemm::threadblock::DefaultMma<ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, LayoutC, arch::OpClassSimt, arch::Sm50, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, false >::ThreadblockMma

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

using cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::OperatorClass = arch::OpClassSimt

Member Data Documentation

template<typename LayoutA , int kAlignmentA, typename LayoutB , int kAlignmentB, typename LayoutC , typename ElementC , typename ArchTag , typename ElementAccumulator , typename ThreadblockShape , typename WarpShape , typename EpilogueOutputOp , typename ThreadblockSwizzle , bool SplitKSerial, typename Operator >

int const cutlass::gemm::kernel::DefaultGemm< int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false >::kEpilogueElementsPerAccess = EpilogueOutputOp::kCount

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