Shared-Memory Tiling (Blocking)

Definition

Shared-memory tiling (a.k.a. blocking) decomposes a matmul into fixed-size output tiles, each computed by one thread block that cooperatively stages sub-tiles of the operands from HBM into on-chip shared memory (SMEM) once, then reuses them across the whole block. It is the rung that converts the naive, zero-reuse GEMM ($I=0.25$ FLOP/byte, deeply memory-bound) into a kernel whose HBM traffic falls by the tile side, marching it rightward across the roofline ridge.

How it works

Each block owns a BM×BN output tile of $C$ and marches the shared $K$ dimension in steps of BK. Per step it stages a BM×BK tile of $A$ and a BK×BN tile of $B$ into SMEM (with coalesced loads), then every thread accumulates its dot-product partials from SMEM. The two-barrier pattern guards the shared buffers:

load As, Bs from HBM        // cooperative, coalesced
__syncthreads()             // (1) all writes done before any read
accumulate As·Bs -> regs
__syncthreads()             // (2) all reads done before next overwrite

Barrier (1) prevents a read-before-write race (fast warp reads garbage); barrier (2) prevents a write-before-read race (next tile clobbers a slow warp’s operands). Both must be reached by all threads — placing one inside a divergent branch deadlocks.

HBM-traffic reduction by tile side. For square tile side $S$, total load traffic is
$\underset{\#\text{tiles}}{\underbrace{\frac{MN}{S^2}}}\cdot \underset{\text{per tile}}{\underbrace{2KS}}=\frac{2MNK}{S},$
versus naive $2MNK$ — a factor-$S$ reduction (each loaded element is reused $S$ times before eviction). Arithmetic intensity rises to
$I_{\text{tiled}}=\frac{2MNK}{(2MNK/S)\cdot b}=\frac{S}{b}\Rightarrow I_{S=32}^{\text{FP32}}=8\frac{\text{FLOP}}{\text{byte}}.$
$S=32$ cuts HBM traffic 32×. But $I=8$ is still below the FP32 ridge (~20–50): the new bottleneck is SMEM bandwidth (2 SMEM loads per FMA), fixed one tier down by register blocking. Bigger tiles reuse more but cost more SMEM (two $S\times S$ FP32 tiles $=2S^2\cdot 4$ B) and registers, hurting occupancy — the central tuning tension. The asymptotic ideal ($I=2n/3b$, hundreds of FLOP/byte) is only approached with the L2 cache capturing reuse across blocks.

Why it matters

Tiling is the foundational move of high-performance GPU compute and the direct ancestor of cuBLAS/CUTLASS GEMM and FlashAttention (which tiles $Q,K,V$ into SRAM so the $N^2$ score matrix never touches HBM). Every transformer training/prefill FLOP flows through tiled GEMM; the difference between a naive and a tiled kernel is ~1% vs. 80%+ of peak — a direct multiplier on GPU-hours.

Taught in

• 05-optimizing-gemm — §5 the tiling scheme, the two-barrier pattern, and the traffic derivation.
• 04-first-cuda-kernels — §9 the naive matmul this fixes.
• 06-memory-wall-and-flashattention — tiling applied to attention.

See also

• register-blocking — the next rung: reuse SMEM values in registers
• gemm — the workload tiling optimizes
• gpu-memory-hierarchy — the SMEM tier tiling exploits
• memory-coalescing — coalesced tile loads
• shared-memory-bank-conflicts — conflict-free SMEM layout
• arithmetic-intensity — the $S/b$ lever
• roofline-model — classifying the tiled kernel
• io-aware-algorithms — the general principle
• gpu-systems-for-llms