metalworkingGitHub ↗
/techniques/tiling

Tiling

Tiling is the reuse pyramid: copy a block of the problem one level closer to the ALUs, do all the math it can support, then let it go. It's the technique that turns bandwidth-bound matmul into compute-bound matmul.

CUDA equivalent: identical concept, and your instincts transfer. What moves is the shape of the pyramid, because Apple's memory hierarchy has different proportions.

The pyramid on this platform, as instantiated by every GEMM in the case studies:

GRID each threadgroup owns a BM×BN output tile device memory (DRAM) one threadgroup THREADGROUP TILE BM×BN = 64×64 · staged slabs of A and B simdgroup simdgroup simdgroup threadgroup memory (32 KB) one simdgroup SIMDGROUP PATCH 32×32 · a 4×4 grid of accumulator fragments 8×8 registers (~208 KB / core) one instruction FRAGMENT simdgroup_float8x8 8×8 multiply-accumulate one MMA per issue each level multiplies reuse: arithmetic per byte of DRAM traffic climbs from O(1) to O(tile edge)

Reading left to right zooms in one level per panel. The strip under each panel names the memory that holds it; the highlighted cell is what the next panel magnifies.

Each level multiplies reuse: a value loaded once into threadgroup memory feeds every simdgroup in the group; a fragment loaded once into registers feeds a whole row or column of accumulator tiles. Arithmetic per byte of DRAM traffic climbs from O(1) (naive) to O(tile edge), which is the entire arithmetic-intensity game.

The Apple-specific calibration, versus CUDA habits:

  • The middle level is thin. 32 KB caps threadgroup tiles; typical shapes are 64×64 output per threadgroup with a 16-32 deep K-slab, smaller than CUDA-typical. The pyramid's weight shifts down a level: register blocking carries more of the reuse than shared-memory blocking does on NVIDIA.
  • Tile shapes are compile-time. Whether via -D defines (m5-gemm) or template parameters (steel's BM/BN/BK/WM/WN), sizes are baked so loops fully unroll and the register allocator can plan. Host code picks the variant per shape.
  • The edge problem is solved at pipeline time, not per-thread. Instead of every thread guarding every access, function constants compile separate aligned/ragged pipelines; the aligned one contains zero bounds checks.

Tiling's two companion moves have their own pages: getting the tile in efficiently is the cooperative load, and hiding the tile's load time behind the previous tile's math is double buffering. The limit case of tiling, where the "tile" you refuse to write to memory is an entire intermediate matrix, is flash attention.

Edit this page on GitHub ↗