Hopper and Blackwell GEMM

TMA fills a ring of shared-memory stages while one thread issues the tensor-core MMAs: warp specialisation, wgmma to tcgen05 and Tensor Memory, and FP8 and FP4 with block scaling.

Check yourself ↓ three questions and what to remember, at the end

Loading the animation…

Concept

Chapter 7 ended with a tensor-core GEMM that every warp runs the same way: load a tile, wait, multiply, repeat. Hopper and Blackwell kernels are built differently, because the hardware gives each job its own engine.

TMA (the Tensor Memory Accelerator, Hopper and later) copies a whole tile from global into shared memory on the request of a single thread, asynchronously, and signals completion on a barrier in shared memory (cp.async.bulk.tensor in PTX). On the A100 every thread of the block took part in a copy (cp.async, 4 to 16 bytes each).
wgmma (Hopper) is issued by a warpgroup of four warps and runs asynchronously; B must come from shared memory, A from registers or shared memory, and the accumulators stay in the warpgroup's registers.
tcgen05.mma (Blackwell) is issued by one thread (PTX ISA Table 55), reads its operands from shared memory (A may be in Tensor Memory) and accumulates in Tensor Memory (TMEM): a dedicated 256 KB per SM, 512 columns by 128 lanes of 32-bit cells, next to the tensor cores. With cta_group::2 one thread of a pair of CTAs on neighbouring SMs starts an MMA that uses both SMs' tensor cores.

So a kernel splits its warps by role, which is called warp specialisation. A producer warp does nothing but issue TMA loads into a ring of shared-memory stages; the MMA issuer waits for a stage to fill, multiplies it and releases it; an epilogue reads finished accumulators, converts them and stores C. Two barriers per stage (full, empty) hand the slots back and forth. The animation runs that loop for one CTA: 3 output tiles of 6 K-slices each.

With the defaults (four stages, two accumulator buffers in TMEM) the tensor core starts as soon as the first load lands and never waits again: busy 78% of the run, the rest being the first load and the last epilogue. Cut the ring to one stage and every load is exposed: 24%. Two stages give 46%, three 64%. Set the accumulators to 1 (registers), the Hopper case with one consumer warpgroup, and each tile's MMAs wait for the previous tile's epilogue: 108 time units instead of 92. Two accumulator buffers let the epilogue drain one while the tensor core fills the other; Hopper kernels get the same effect with two consumer warpgroups taking turns ("ping-pong"), Blackwell's TMEM holds both accumulators without spending registers. With a long epilogue (set it to 30, longer than a tile's 6 slices of MMA) even two buffers expose some of it: 126 units, against 174 with one.

Concept

How many stages fit. A stage holds one K-slice of A and of B. For a 128 × 256 output tile with 64-deep BF16 slices that is 48 KB, so the B200's 228 KB of shared memory per SM (compute capability 10.0) holds 4 stages, the H100's 4 and the A100's 3. The steady-state equation says how many are needed: enough that $S$ slices of MMA cover one load's latency plus its MMA. Below that the ring runs dry and the tensor core idles (the hatched gaps); above it, extra stages buy nothing. On Blackwell the accumulators no longer compete with the stages for space: they live in TMEM, and the registers are left for the epilogue.

The ridge moved again. The tensor cores' peak divided by HBM bandwidth is the arithmetic intensity a GEMM needs before it stops waiting on memory (chapter 2). It has risen every generation, and each narrower format doubles it:

GPU	BF16 dense	FP8 dense	NVFP4 dense	HBM	Ridge, BF16	Ridge, FP8	Ridge, NVFP4
A100	312 TFLOP/s	none	none	1.555 TB/s	201
H100	990 TFLOP/s	1,980 TFLOP/s	none	3.35 TB/s	295	591
B200 (GB200)	2,500 TFLOP/s	5,000 TFLOP/s	10,000 TFLOP/s	8 TB/s	313	625	1,250

Ridges are in flops per byte. The B200 is the GPU of a GB200 NVL72, with the same figures the AI Racks and AI Accelerators sites use; the air-cooled HGX B200 board is rated lower. Its BF16 peak is 2.53 times the H100's, and NVFP4 is 4 times BF16 again. A kernel that reached the H100's ridge has to find more reuse to reach the B200's: bigger tiles, CTA pairs that share operands, L2 reuse across the blocks of a wave.

Concept

FP8 and FP4, with block scaling. An FP8 or FP4 value has too little range to hold a tensor on its own, so it carries a scale. Hopper's FP8 uses one scale per tensor, kept by software. Blackwell's tensor cores apply block scales inside the MMA (tcgen05.mma with .kind::mxf8f6f4, .kind::mxf4 or .kind::mxf4nvf4): every row of A and column of B is cut into blocks, and each block's values are multiplied by its own scale.

Format	Element bits	Elements per scale	Scale	Bits per value	An 8192 × 8192 matrix
BF16	16		none	16	134 MB
FP8 E4M3, per-tensor scale	8	whole tensor	FP32	8	67.1 MB
MXFP8 (OCP)	8	32	E8M0 (a power of two)	8.25	69.2 MB
MXFP4 (OCP)	4 (E2M1)	32	E8M0	4.25	35.7 MB
NVFP4	4 (E2M1)	16	E4M3, plus one FP32 per tensor	4.5	37.7 MB

NVFP4's smaller blocks and finer E4M3 scale track the data more closely than MXFP4's power-of-two scale every 32 values, for a quarter of a bit more per value. What these formats do to accuracy is Numerics Explained, chapter 10; chapter 11 here is what fewer bytes do to a decode step.

Maths

Number the K-slices $g = t K + k$ (tile $t$ , slice $k$ ), with $S$ stages and $A$ accumulator buffers. The schedule is a set of max-plus recurrences:

\begin{aligned} L_g &= \max(L_{g-1} + t_{\text{issue}},\ D_{g-S}) && \text{(a slot is free once its MMA is done)}\\ M_g &= \max(L_g + t_{\text{load}},\ D_{g-1},\ E_{t-A} \text{ if } k = 0) && \text{(data in, tensor core free, buffer drained)}\\ D_g &= M_g + t_{\text{mma}}, \qquad E_t = \max(D_{tK+K-1},\ E_{t-1}) + t_{\text{epi}}. \end{aligned}

Within a long tile, chase one slot round the ring: slice $g + S$ cannot load before slice $g$ 's MMA is done, and then needs $t_{\text{load}} + t_{\text{mma}}$ more, so $S$ consecutive slices take at least $t_{\text{load}} + t_{\text{mma}}$ ; they also take at least $S\,t_{\text{mma}}$ (one tensor core) and $S\,t_{\text{issue}}$ (one TMA issue stream). The tests check that the model settles to exactly $M_{g+S} - M_g = \max(S\,t_{\text{mma}}, S\,t_{\text{issue}}, t_{\text{load}} + t_{\text{mma}})$ , which gives the equation above. The tensor core is never starved once $S \ge (t_{\text{load}} + t_{\text{mma}})/t_{\text{mma}}$ : here $(12 + 4)/4 = 4$ stages. With $A = 2$ and $t_{\text{epi}} \le K\Delta$ , the epilogue of tile $t$ finishes before tile $t + 2$ needs its buffer, and the run is just $t_{\text{load}} + T K\, t_{\text{mma}} + t_{\text{epi}}$ , which a test also checks.

Code

The recurrences, cut from src/lib/gpu/blackwell.ts (a line-for-line port of reference/blackwell.py):

let a = g > 0 ? ls[g - 1]! + issue : 0;
if (g >= stages && md[g - stages]! > a) a = md[g - stages]!;
ls.push(a);
ld.push(a + load);
let b = ld[g]!;
if (g > 0 && md[g - 1]! > b) b = md[g - 1]!;
if (k === 0 && t >= accs && ed[t - accs]! > b) b = ed[t - accs]!;

The shape of a Blackwell main loop, illustrative and not compiled by this site's CI (the names follow CUTLASS's pipeline classes; the real calls take descriptors and barrier phases):

if (warp == PRODUCER && lane == 0) {          // one thread issues every load
  for (int g = 0; g < slices; ++g) {
    empty[g % S].wait();                      // the slot's last MMA is done
    tma_load(smemA[g % S], smemB[g % S], full[g % S]);  // completes on full[]
  }
} else if (warp == MMA && lane == 0) {        // one thread issues every MMA
  for (int g = 0; g < slices; ++g) {
    full[g % S].wait();
    tcgen05_mma(tmem_acc[tile % 2], smemA[g % S], smemB[g % S]);
    tcgen05_commit(empty[g % S]);             // free the slot when it finishes
  }
}  // epilogue warpgroup: tcgen05.ld from TMEM, convert, store C

Go deeper

PTX ISA: §9.7.17 wgmma, §9.7.18 the 5th-generation TensorCore family (tcgen05, Tensor Memory, Table 55's issue granularity) and cp.async.bulk.tensor (TMA); the CUDA Programming Guide on asynchronous copies.
CUTLASS, Efficient GEMM in CUDA: pipelining, warp specialisation and the ping-pong and cooperative kernel schedules.
NVIDIA, Inside Blackwell Ultra (Table 2: the Blackwell and Blackwell Ultra peaks; NVFP4's two-level scaling) and the GB200 NVL72 specifications; Jarmusch and Chandrasekaran, Microbenchmarking NVIDIA's Blackwell Architecture (2025: 148 SMs on the B200, TMEM and tcgen05 measured); Rouhani et al., Microscaling Data Formats for Deep Learning (2023: the OCP MX formats).
NVIDIA GPU series: 5th-gen tensor cores and MX formats, the Blackwell SM, TMA and MXFP4 against NVFP4.
Systolic Arrays Explained, chapter 8: tensor cores against systolic arrays; a single-thread MMA into dedicated accumulator memory is a step towards the systolic side.

Check yourself

Commit to an answer before you reveal it: a wrong guess you can explain teaches more than a right one you skimmed. Every number in the answers comes from the same tested model as the chapter’s.

Q1Cut the ring to one stage. For what share of the run is the tensor core busy?

Your answermove the slider

Reveal the answer

Answer: 24% (counted right within ± 4%)

24%: with one stage every load is exposed. The default four stages keep it busy 78% of the run, the rest being the first load and the last epilogue.

The pipeline's defaults →

Commit to a number, then reveal.

Q2How many stages of a 128 × 256 output tile with 64-deep BF16 slices fit in a B200 SM's shared memory?

Your answer

Reveal the answer

Answer: 4 (exact)

4: one stage is 48 KB and the SM has 228 KB. The H100 also holds 4 and the A100 3.

How many stages fit →

Commit to a number, then reveal.

Q3Use one accumulator buffer in registers (the Hopper case with one consumer warpgroup) instead of two in Tensor Memory. What happens?

The run is just as fast
Each tile's MMAs wait for the previous tile's epilogue, so the run is longer
The tensor core never waits

Reveal the answer

Answer: Each tile's MMAs wait for the previous tile's epilogue, so the run is longer

The run is longer: 108 time units against 92. A second buffer lets the epilogue drain one accumulator while the tensor core fills the other.

Warp specialisation →

Pick one, then reveal.