Understanding cutlass GEMM hierarchy

uniadam · October 14, 2021, 4:58pm

I think this picture is showing what cutlass is doing.

But I am not understanding what is happening. Or what is the shape?

Here they are defining several shape, why several and how it is going to work?

NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_fprop_implicit_gemm_f16nhwc_f16nhwc_f16nhwc_tensor_op_f16_sm80.cu#L290


      
          
          
/// Conv operation element types for the Gemm equivalent (ImplicitGemm)
          using ElementA           = cutlass::half_t;
          using ElementB           = cutlass::half_t;
          using ElementC           = cutlass::half_t;
          using ElementAccumulator = cutlass::half_t;
          using ElementCompute     = cutlass::half_t;
          
          
/// Device-level Conv2d instance
          using Conv2dFpropKernel = typename cutlass::conv::kernel::DefaultConv2dFprop<
            ElementA, cutlass::layout::TensorNHWC,
            ElementB, cutlass::layout::TensorNHWC,
            ElementC, cutlass::layout::TensorNHWC,
            ElementAccumulator,
            cutlass::arch::OpClassTensorOp,
            cutlass::arch::Sm80,
            cutlass::gemm::GemmShape<128, 128, 64>,
            cutlass::gemm::GemmShape<64, 64, 64>,
            cutlass::gemm::GemmShape<16, 8, 16>,
            cutlass::epilogue::thread::LinearCombination<
              ElementC,

cutlass::gemm::GemmShape<128, 128, 64>,
cutlass::gemm::GemmShape<64, 64, 64>,
cutlass::gemm::GemmShape<16, 8, 16>,

Or here how can I relate each line to the picuture that I mentioned?

ShapeMMAThreadBlock
ShapeMMAWarp 
ShapeMMAOp

github.com

NVIDIA/cutlass/blob/master/examples/07_volta_tensorop_gemm/volta_tensorop_gemm.cu#L150


      
          using ElementInputA = cutlass::half_t;              // <- data type of elements in input matrix A
          using ElementInputB = cutlass::half_t;              // <- data type of elements in input matrix B
          using ElementOutput = float;                        // <- data type of elements in output matrix D
          
          
// The code section below describes matrix layout of input and output matrices. Column Major for
          // Matrix A, Row Major for Matrix B and Row Major for Matrix C
          using LayoutInputA = cutlass::layout::ColumnMajor;
          using LayoutInputB = cutlass::layout::RowMajor;
          using LayoutOutput = cutlass::layout::RowMajor;
          
          
// This code section describes whether you want to use tensor cores or regular SIMT cores on GPU SM
          using MMAOp = cutlass::arch::OpClassTensorOp;
          
          
// This code section describes CUDA SM architecture number
          using SmArch = cutlass::arch::Sm70;
          
          
// This code section describes the tile size a thread block will compute
          using ShapeMMAThreadBlock =
              cutlass::gemm::GemmShape<128, 128, 32>;  // <- threadblock tile M = 128, N = 128, K = 32
          // This code section describes tile size a warp will compute
          using ShapeMMAWarp = cutlass::gemm::GemmShape<64, 64, 32>;  // <- warp tile M = 64, N = 64, K = 32

mnicely · October 14, 2021, 5:04pm

Maybe these will provide some insight.

https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/