Register-blocked matmul

Two tiles per thread

With 256 threads arranged 16 by 16 and TM = TN = 8, each thread computes an 8 by 8 sub-tile and the block covers 128 by 128. The accumulators are the kernel:

float acc[8][8] = {};                       // 64 registers before addresses
for (int k0 = 0; k0 < N; k0 += 8) {
  // 256 threads load As[8][128] and Bs[8][128], 16 B per thread per tile
  __syncthreads();
  #pragma unroll
  for (int k = 0; k < 8; ++k) {
    float a[8], b[8];                       // filled from As[k][...] and Bs[k][...]
    #pragma unroll
    for (int i = 0; i < 8; ++i)
      for (int j = 0; j < 8; ++j) acc[i][j] = fmaf(a[i], b[j], acc[i][j]);
  }
  __syncthreads();
}

Why intensity rises

Per k-step the thread reads 16 shared values and issues 64 FMAs: four per read, against two reads per FMA for a one-output thread. With the block size fixed, the block tile grows with the per-thread tile, so global traffic per FLOP falls roughly with the tile width rather than its area, and the shared reads per FMA fall from two to (RM + RN) / (RM * RN). The 64 accumulators are also the instruction-level parallelism: each acc[i][j] is a chain the scheduler can interleave.

The bill arrives in registers

-Xptxas -v prints registers per thread and any spill stores. Sixty-four accumulators plus fragments plus addresses is how a kernel reaches the register ceiling and falls to one block per SM. __launch_bounds__(256, 2) makes ptxas fit two blocks and pays with spills if the tile is too large for the budget; that trade is gpu-10.

The flags above are stable, but register counts are per-architecture and per-release, so re-run -Xptxas -v with the toolkit you are actually using.

Things that break

An acc array indexed by a loop the compiler cannot unroll becomes local memory, which lives in DRAM, and -Xptxas -v reports a stack frame and local loads when it happens. A tile that is too large does the same by exhausting registers. Neither failure is visible in the source; both are visible in the ptxas output. And none of this guarantees cuBLAS parity: it moves the kernel from the bandwidth roof toward the compute roof, where gpu-07 measures what remains.