✳flâneur — a map of the web's best reading

Implementing a fast Tensor Core matmul on the Ada Architecture | spatters.ca

spatters.ca · 5,826 words · saved by 1 readers

Using Tensor Cores is now a prerequisite to get anywhere near peak performance on NVIDIA GPUs. In this post we work through the process of developing an efficient Tensor Core matrix multiplication kernel targeting the Ada architecture. We start with a naive implementation and by incorporating techniques used in CUTLASS1, finish with a kernel that matches cuBLAS performance (on one particular problem specification): In the process we’ll learn about the mma, ldmatrix and cp.async PTX instructions, how CUTLASS’s permuted shared memory layout avoids bank conflicts and how to set up an n-stage global to shared memory pipeline. The code is written as simply as possible: the aim is ease of understanding rather than generality or robustness. As may be clear already, this post was heavily inspired by Simon Boehm’s great worklog on optimizing a CUDA matmul kernel2. We’ll focus on one particular problem shape: M=N=K=4096, for fp16 A/B and fp32 C/D. This operation is 2*4096^3 = 137.4 GFLOP3 (conve

Using Tensor Cores is now a prerequisite to get anywhere near peak performance on NVIDIA GPUs. In this post we work through the process of developing an efficient Tensor Core matrix multiplication kernel targeting the Ada architecture. We start with a naive implementation and by incorporating techniques used in CUTLASS 1 , finish with a kernel that matches cuBLAS performance (on one particular problem specification): Kernel Execution Time TFLOP/s % cuBLAS % 4090 peak cublasGemmEx 895 us 153.6 100% 93.0% Kernel 1.0: Naive mma 4680 us 29.4 19.1% 17.8% Kernel 1.1: Naive + 2x tiling 2400 us 57.3 3

Explore this link on the map →

saved by

related reading