Part II · the five hardware primitives

Tensor cores across generations — mma → wgmma → tcgen05

The tensor-core operation D = A·B + C has read the same since Volta. Yet a kernel that saturates one generation is slow — or silently wrong — on the next. The reason is that the tensor core consumes operands in very specific hardware layouts, and three things kept moving underneath that one equation: how the MMA is issued, how the operands are laid out, and where the accumulator lives. Following that drift is what forces every later primitive in this track.

Two constraints that never went away

Before anything moves, fix the two costs that are invariant across all three generations — they are the reason layout is non-negotiable in the first place:

• Global-memory coalescing. A warp's 32 threads should touch one contiguous, aligned HBM segment per access, or the memory system replays the transaction and effective bandwidth collapses. True on Volta, true on Blackwell.
• SMEM bank conflicts. Shared memory is 32 banks of 4 bytes. If two lanes in a warp hit different addresses in the same bank, the access serializes. The fix is the swizzle from lesson 02: physical_col = logical_col ⊕ row, which scatters a column across all 32 banks so the tensor core reads it conflict-free.

Every generation pays these two taxes. What changes is the machinery around the MMA that decides who pays them and when.

The invariant and the three moving parts

The math never moves: a tile of A (M×K) times a tile of B (K×N) accumulates into C (M×N), producing D. The hardware contract around it has three knobs, and each generation turns one:

Read the table top to bottom and one column jumps out: the accumulator marches out of the register file — regs → regs → TMEM. That is the through-line of this lesson, and the widget below is built to make it visible. But to see why each move was forced, walk the rows one at a time.

Generations · where the operands and accumulator live

Pick a generation. The diagram redraws where A, B, and the accumulator physically sit (register file vs an SMEM descriptor vs TMEM), and which agent commits the MMA (a full warp vs a 128-thread warpgroup vs a single elected thread). The accumulator box is highlighted as it migrates out of registers — the one change that forces lesson 05.

Ampere — registers, fragments, and ldmatrix

On Ampere the whole MMA lives in the register file. The instruction is mma.sync.aligned.m16n8k16; the .sync means all 32 lanes of the warp execute it together, synchronously — it returns when the math is done. (This is the world the gpu_kernels tensor-core lesson teaches; we extend it forward from here.)

Because the operands are register fragments, each lane owns a precise slice of each matrix. For the m16n8 fp32 accumulator the tile is 16×8 distributed over 32 lanes, 4 fp32 values per lane. The exact map: lane l holds rows { l/4 , l/4 + 8 } and columns { 2·(l mod 4) , 2·(l mod 4)+1 }. Each 8×8 half of that fragment — the m8n8 tile that recurs on every generation — is, in CuTe layout notation,

S[(8,4,2):(4@laneid, 1@laneid, 1@m)]

You do not place operands into that layout by hand. The PTX ldmatrix instruction does it:

// load A/B fragments SMEM → registers, doing the cross-lane shuffle
ldmatrix.sync.aligned.m8n8.x1.shared.b16   {r0},               [addr];
ldmatrix.sync.aligned.m8n8.x2.shared.b16   {r0,r1},            [addr];
ldmatrix.sync.aligned.m8n8.x4.shared.b16   {r0,r1,r2,r3},      [addr];
// optional .trans transposes the 8×8 tile as it lands (for B / col-major)
ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {r0,r1,r2,r3}, [addr];

ldmatrix loads one, two, or four 8×8 tiles of 16-bit elements and shuffles them across lanes into the fragment registers mma.sync expects. The addressing is the subtle part: for matrix m, row r, the base address is supplied by lane m·8 + r — so .x1 reads its 8 row-base pointers from lanes 0–7, .x2 from lanes 0–15, and .x4 from all of lanes 0–31. You give each lane the SMEM address of one row; the hardware does the transpose-and-scatter.

Two consequences set up everything that follows:

• Swizzle is hand-managed. There is no descriptor naming the SMEM layout — the kernel computes swizzled addresses with index math (the ⊕ row from lesson 02) before handing them to ldmatrix. Get the math wrong and you read conflicted, or simply wrong, bytes.
• There is no reverse of ldmatrix. To write the result out, the kernel stores the accumulator with ordinary per-thread stores — each lane writes its 4 fp32 values to the right global addresses. (Hold onto this 8×8 fragment shape; it returns in every generation.)

Hopper — the SMEM matrix descriptor, and swizzle becomes a name

Ampere's ldmatrix step is pure overhead: it occupies the warp moving bytes that the tensor core is about to consume anyway. Hopper removes it. wgmma ("warpgroup MMA") reads its operands directly from shared memory through a matrix descriptor — a 64-bit handle the kernel builds once that tells the hardware how the SMEM tile is laid out:

• start address (in SMEM),
• leading-dimension byte offset,
• stride-dimension byte offset,
• swizzle mode,
• base offset.

An mma is now issued by a whole warpgroup — 128 threads (4 warps) acting as one — and it is asynchronous: it returns immediately and completes later, which is exactly why lesson 06's mbarriers exist. The two operand-source variants:

• wgmma.mma_async.sync.aligned.m64nNk16....ss — both A and B come from SMEM descriptors.
• ....rs — A comes from registers, B from an SMEM descriptor (useful when A is freshly produced and already resident in registers).

The accumulator is still a register fragment — that is the one thing Hopper does not change, and it is the wall Blackwell will hit. The shape family is wgmma.mma_async.sync.aligned.m64nNk16 with N ∈ {8 … 256}, so one instruction can produce up to a 64×256 tile.

The decisive new rule is what the descriptor's swizzle mode implies. On Ampere swizzle was loose index math; on Hopper it becomes a named format:

Blackwell — tcgen05.mma, one issuing thread, and the accumulator leaves registers

Hopper kept the accumulator in registers, and that became the ceiling. As MMA tiles grew, a register-held accumulator ate the register budget, which throttled occupancy and capped tile size. Blackwell breaks the link: tcgen05.mma writes its accumulator to TMEM — a dedicated tensor-core memory that is not the register file (lesson 05 is entirely about it). This is the big change.

Three more shifts ride along:

• One elected thread issues the MMA on behalf of the group — not a warp, not a warpgroup. (Electing exactly one lane, and committing exactly once, becomes a recurring correctness trap in lesson 13.)
• It can span two CTAs. cta_group::1 is a single CTA; cta_group::2 lets two cooperating CTAs in a cluster jointly produce one large tile — the seed of lesson 07's distributed shared memory.
• Fully asynchronous, like wgmma. A and B are staged into SMEM (swizzled, brought in by TMA) and read via descriptors, exactly as on Hopper.

Since the accumulator now lives in TMEM, you need an instruction to get it back for the epilogue:

// accumulator stays in TMEM during the K-loop, then:
tcgen05.ld.sync.aligned...   // TMEM → registers  (epilogue read-back)
tcgen05.st.sync.aligned...   // registers → TMEM  (e.g. correction write-back)
tcgen05.cp...                // SMEM → TMEM        (stage block-scale factors)

Block-scaled low precision

Blackwell also adds block-scaled narrow formats, where a small group of elements shares one scale factor:

• MXFP8 — fp8 values with an e8m0 (power-of-two) scale per block.
• NVFP4 — fp4 values with an e4m3 scale per block.

This adds two new operands to the MMA: scale tensors SFA(M, K/B) and SFB(N, K/B) for block size B. They are staged SMEM → TMEM via tcgen05.cp, and a scale_vec selector packs them at 1X / 2X / 4X granularity. The point for this lesson: the scale factors are more operands with their own layout contract — the descriptor discipline does not get simpler at low precision, it gets one tensor wider.

The recurring m8n8 fragment

One small object survives all three generations: the 8×8 register fragment. Trace it:

That continuity is not trivia — it is why the same cast-and-store epilogue code (downcast fp32 → bf16, apply bias/activation, write to global) is reusable across all three generations. The accumulator's home changed; the shape it arrives in for the epilogue did not. Lane l still lands on row l/4 of that fragment after a tcgen05.ld, exactly as it did after mma.sync on Ampere.