GEMM: CTA Tiles, Register Tiles, and Tensor Cores

Core Concept

High-performance GEMM uses hierarchical tiling: a CTA owns an output tile, warps and threads hold smaller accumulator tiles, and the K dimension advances through staged tiles.

Why It Affects Performance

Performance comes from reusing A and B while issuing many FMAs. Layout, tile shape, pipeline stages, and accumulator count determine Tensor Core utilization and resource pressure.

Example or Pseudocode

acc = tl.zeros((BM, BN), tl.float32)
for k in range(0, K, BK):
    a = tl.load(a_ptrs); b = tl.load(b_ptrs)
    acc = tl.dot(a, b, acc)
tl.store(c_ptrs, acc)

Profiling and Performance Judgment

Use cuBLAS as a strong baseline. Inspect HMMA/Tensor instructions, Tensor Pipe utilization, DRAM bytes, occupancy, and achieved TFLOP/s. Build sm_75 code for the RTX 2070 SUPER.

Five-Minute Exercise

Why are FP16 inputs with FP32 accumulators common?