default_gemm.h File Reference

Default kernel-level GEMM definitions combine threadblock-scoped matrix multiply-add with the appropriate threadblock-scoped epilogue. More...

`#include "cutlass/cutlass.h"`
`#include "cutlass/layout/matrix.h"`
`#include "cutlass/numeric_types.h"`
`#include "cutlass/arch/wmma.h"`
`#include "cutlass/epilogue/threadblock/epilogue.h"`
`#include "cutlass/epilogue/thread/linear_combination.h"`
`#include "cutlass/gemm/gemm.h"`
`#include "cutlass/gemm/kernel/gemm.h"`
`#include "cutlass/gemm/kernel/gemm_pipelined.h"`
`#include "cutlass/gemm/threadblock/default_mma_core_sm75.h"`
`#include "cutlass/gemm/threadblock/default_mma_core_sm70.h"`
`#include "cutlass/gemm/threadblock/default_mma.h"`
`#include "cutlass/gemm/threadblock/default_mma_core_simt.h"`
`#include "cutlass/gemm/threadblock/threadblock_swizzle.h"`
`#include "cutlass/epilogue/threadblock/default_epilogue_tensor_op.h"`
`#include "cutlass/epilogue/threadblock/default_epilogue_volta_tensor_op.h"`
`#include "cutlass/epilogue/threadblock/default_epilogue_simt.h"`
`#include "cutlass/transform/threadblock/predicated_tile_iterator.h"`
Include dependency graph for default_gemm.h:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

## Classes

struct | cutlass::gemm::kernel::DefaultGemm< ElementA_, LayoutA_, kAlignmentA, ElementB_, LayoutB_, kAlignmentB, ElementC_, LayoutC_, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, SplitKSerial, Operator, IsBetaZero >
struct | cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >
| Partial specialization for Turing Architecture. More...

struct

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 >

| Partial specialization for Turing Integer Matrix Multiply Interleaved layout. More...
struct | cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm70, ThreadblockShape, WarpShape, GemmShape< 8, 8, 4 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >
| Partial specialization for Volta architecture. More...
struct | cutlass::gemm::kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape< 1, 1, 1 >, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator >
| Partial specialization for SIMT. More...
struct | 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 >
| Partial specialization for SIMT DP4A. More...

## Namespaces

| cutlass
| cutlass::gemm
| cutlass::gemm::kernel

## Detailed Description

Note, CUTLASS epilogues universally target row-major outputs. Column-major outputs are accommodated by exchanging A and B operands and assuming transposed layouts. Partial specializations here choose 'device::GemmTransposed' to implement this functionality.