Case Study: The Double-Buffered GEMM

sync_copy_db.metal is the tiled GEMM with double buffering added, and the cleanest measured demonstration in this glossary that there is no single fastest kernel.

The delta from its sibling is exactly three structural changes: two buffer sets, a prologue prefetch, one barrier per iteration instead of two.

  threadgroup float A_tg[2][BM * BK];
  threadgroup float B_tg[2][BK * BN];
  ...
  // Prologue: load tile 0 into buffer 0.
  load_tile<BM, BK, NTHREADS>(A + tg_row * k, k, A_tg[0], tid);
  load_tile<BK, BN, NTHREADS>(B + tg_col,     m, B_tg[0], tid);

  for (uint l = 0; l < k_tiles; l++) {
    threadgroup_barrier(mem_flags::mem_threadgroup);
    ushort nxt = ushort(1) - cur;
    if (l + 1 < k_tiles) {
      load_tile<BM, BK, NTHREADS>(A + tg_row * k + k_off, k, A_tg[nxt], tid);
      load_tile<BK, BN, NTHREADS>(B + k_off * m + tg_col, m, B_tg[nxt], tid);
    }
    /* compute on A_tg[cur] / B_tg[cur] */
    cur = nxt;
  }

— sync_copy_db.metal:87-126↗, abridged

The single barrier is doing double duty: "the prefetch landed" and "everyone's done with the buffer we're about to overwrite" are the same condition when the buffers alternate. The kernel's own header comment is honest about the mechanism: without the dead async-DMA primitive this is not true overlap. It's ILP-based latency hiding, the same threads issuing independent loads and math for the scheduler to interleave.

The payoff table (author's numbers, M5 Max, fp32, best of 3×5, from the repo README↗):

Three regimes, three winners. At 1024², launch and DRAM-fetch latency dominate and the prefetch pipeline hides them: 3× over both alternatives, and note this is the regime LLM decode lives in. At 4096², the simpler kernel wins because its smaller loop unrolls better; compiler quality beats algorithm. At 8192², everything hits the bandwidth wall and MPS converges with the hand-written kernels.

The README's "things that did not help" section is required reading as methodology: offline -O3 -ffast-math compilation (bit-identical to runtime), bigger simdgroup tiles (10× slower from spills), SW=3 threadgroups. Negative results, measured and published: the norm this glossary's war stories wish the whole ecosystem followed.