# Benchmarking GPU Tensor Cores on General Matrix Multiplication Kernels through CUTLASS
## 2 Materials and Methods
### GPU Architecture and Programming Model
A typical procedure of producing the GEMM results of cutlass::half_t (i.e., fp16) as input type is presented below:
```
usingElementA=cutlass::half_t; usingLayoutA=cutlass::layout::RowMajor; constexprintAlignmentA=128/cutlass::sizeof_bits<ElementA>:: value; //SameasmatricesBandC usingElementAccumulator=cutlass::half_t; usingArchTag=cutlass::arch::Sm80; usingOperatorClass=cutlass::arch::OpClassTensorOp; usingThreadblockShape=cutlass::gemm::GemmShape<128,128,32>; usingWarpShape=cutlass::gemm::GemmShape<64,64,32>; usingInstructionShape=cutlass::gemm::GemmShape<16,8,16>; constexprintNumStages=4; usingEpilogueOp= cutlass::epilogue::thread::LinearCombination< ElementC,AlignmentC,ElementAccumulator,ElementAccumulator>; //Classicdata-paralleldeviceGEMMimplementationtype usingDeviceGemmBasic= cutlass::gemm::device::GemmUniversal<ElementA,LayoutA,ElementB, LayoutB,ElementC,LayoutC,ElementAccumulator,OperatorClass,ArchTag,ThreadblockShape,WarpShape,InstructionShape, EpilogueOp, cutlass::gemm::threadblock:: GemmIdentityThreadblockSwizzle<>,NumStages,AlignmentA,AlignmentB>;
```

As we claimed before, CUTLASS is a highly templated library; hence, nearly all of its data types and interfaces are specified with templates. The matrix is defined by the tuple <Element, Layout> and AlignmentA, and it is used to determine the memory access granularity to the corresponding memory region. Since CUDA now supports 128-bit coalesced memory access (using LDG.128 SASS instruction), the number of elements per instruction loads is thus determined by 128/sizeof_bits(Element). Next, some required information is defined such as the MMA operation type (whether CUDA Core or Tensor Core will be used) and the underlying GPU hardware architecture tag. Lastly, the hierarchical tile shapes assigned to different levels of the collaboration group (thread block, warp or per MMA) are defined with the GemmShape template type. The epilogue operation is the auxiliary phase after the main loop of kernel for the scalar calculation or type conversion. Note that asynchronous global memory to shared memory instruction (LDG.STS SASS instruction) is proposed since Ampere generation; therefore, this producer-consumer pattern is pipelined and several memory access stages are exploited to overlap \begin{table}
\begin{tabular}{l l} \hline \hline
**API Level** & **API Class/Function** \\ \hline Device & device::GemmUniversalAdapter \\ \hline Kernel & kernel::GemmUniversal \\ \hline Collective & collective::CollectiveMma \\ \hline Tile (MMA and copy) & cute::TiledMma,cute::TiledCopy \\ \hline Atom & cute::Mma_Atom,cute::Copy_Atom \\ \hline \hline \end{tabular}
\end{table} 
Table 1: CUTLASS API overview. the MMA operation and global memory access. The number of pipelined stages is defined by NumStages. All the above information is then passed to GEMM kernel templates to generate the target code. There are several advantages of encoding key information (such as element types, tile shapes and memory access granularities) as template parameters. For example, if the tile shape is fixed at the compile-time, register spills could be avoided when indexing tile elements. Besides, loop unrolling could be performed accordingly to saturate the instruction cache. Further, with this compiler-assistant information, register usage could be reduced since some variables are encoded as constants and part of input data could be fetched from CUDA constant memory to mitigate the DRAM bandwidth. ## 3 Results and Discussion

In this section, we demonstrate the key insights discovered from the benchmarks. ### Testbed

Our experiments are conducted with NVIDIA A100 40 GB PCIe version. There are 108 Tensor Core SMs and 40 MB L2 cache total, with 1555 GB/s off-chip memory bandwidth. On the software side, we built CUTLASS 3.0 with host compiler [EMAIL_ADDRESS] and CUDA version '11.4.0' under Debian 11 as the host OS. All measured kernels were warmed up with 10 repeated runs and invoked 20 times to obtain the average execution elapsed times.