Part II · the five hardware primitives

mbarriers and the phase model

An asynchronous engine that returns immediately needs an asynchronous way to tell you it has finished. The classic block-wide __syncthreads() cannot do this job — it is thread-only, stateless, and can name only a single rendezvous point. The replacement is the mbarrier: a hardware sync object that counts arrivals, counts bytes, and carries a phase bit so it can be safely reused round after round. That phase bit is the single most important — and most error-prone — idea in the whole async paradigm. Get it wrong and the kernel computes silent garbage.

Why __syncthreads() is the wrong tool

The barrier from the gpu_kernels world is a synchronous, all-threads-in-the-block, single-point rendezvous. Line up every thread, release them together, done. For the async supply chain it fails on four counts at once:

An mbarrier ("memory barrier," sometimes "async barrier") is a 64-bit object that lives in shared memory. The kernel allocates it, initializes it, and from then on producers arrive on it and consumers wait on it. Its entire state is two numbers packed into those 64 bits.

The two pieces of state

Everything about correctness flows from understanding exactly what an mbarrier holds:

• An arrival counter — how many arrivals are still missing this round. init(count) sets the expected arrival count; each arrive decrements it.
• A phase bit — which round the barrier is in: 0, then 1, then 0, … It starts at phase 0 after init.

In the toolchain the object is set up with a single PTX instruction, and the count is the number of arrivals that constitute "done" for one round:

// allocate a 64-bit mbarrier in shared memory, expect `count` arrivals, start in phase 0
mbarrier.init.shared.b64 [bar], count;

The operations

Three operations act on the object. The second one does double duty and is the source of most confusion, so read it twice:

• arrive — a producer signals arrival, decrementing the arrival counter by one.
• arrive.expect_tx(bytes) — does two jobs in one instruction: it records the expected transaction byte count for this round and counts as the issuing thread's arrival. The round is complete only when both conditions hold: arrival count → 0 and pending byte count → 0. As the TMA engine streams the tile in, it issues complete-tx updates against the same barrier that decrement the pending byte count — so the engine, not a thread, drives the byte half to zero.
• wait — try_wait(bar, phase) blocks until the barrier's internal phase differs from the phase argument. Read that carefully: the argument names the phase you expect to leave behind, not the one you are reaching. You pass the phase you are currently in; the call returns once the round completes and the bit flips away from it.

The phase flip rule

Here is the mechanism that makes a single barrier reusable forever. Each time a barrier completes all arrivals (and all bytes) for its round, it flips:

phase: 0 → 1 → 0 → 1 → … (toggle on every completed round)

The kernel does not read the bit back from the barrier. Instead it keeps the expected phase in a register and toggles that register after each successful wait. First use is phase 0, so:

• iteration 0 waits for phase 0 (it expects to leave phase 0 behind),
• iteration 1 waits for phase 1,
• iteration 2 waits for phase 0 again, and so on.

In code this is one line — the line whose omission is the canonical bug of the entire paradigm:

int phase = 0;
for (int it = 0; it < num_iters; ++it) {
    // ... producer arrives on `bar` for this round ...
    mbarrier.try_wait.parity.shared.b64  done, [bar], phase;   // block until bit != phase
    // ... consume the freshly-ready data ...
    phase ^= 1;                                                // TOGGLE — distinguish next round
}

The PTX spin is the non-blocking mbarrier.try_wait.parity.shared.b64 in a loop — "parity" is NVIDIA's name for the phase bit. CUTLASS wraps this whole dance as PipelineTmaAsync so you rarely write the spin by hand, but the phase register is still yours to manage.

Interactive · the mbarrier state inspector and the phase-reuse bug

This is the centerpiece of the lesson. Drive an mbarrier by hand: init it, have a thread arrive.expect_tx to announce a byte budget, let the TMA engine deliver bytes in chunks, and wait. Watch the arrival counter and pending-byte count fall to zero, the phase bit flip, and the round complete. Then step through several pipeline iterations and watch the bit toggle 0 → 1 → 0 …

Now flip the "forget phase ^= 1" toggle. The consumer stops toggling its expected phase, so on iteration 1 it waits for phase 0 — which the barrier already left behind in iteration 0. The wait returns immediately on last round's completion, before the new load has landed. The inspector flags the early return: the consumer reads a half-filled tile. That is silent corruption.

Run it once cleanly, then again with the bug toggle on and watch iteration 1's wait return on the stale completion — the byte bar is still red (data not yet delivered) when the consumer is released.

tcgen05 does not arrive on its own

The TMA engine politely issues complete-tx updates as it delivers bytes, so a load self-completes its barrier. The tensor core does not. A tcgen05 MMA finishes asynchronously but advances no barrier unless you explicitly attach an arrival to its commit path:

// the MMA completes async — you must route its completion to a barrier:
tcgen05.commit.mbarrier::arrive  [mma_done];   // the MMA path arrives here when math is done
// the epilogue then waits on mma_done before reading TMEM
// (tcgen05.cp — SMEM→TMEM scale-factor copies — attaches an arrival the same way)

So the rule splits by agent: TMA loads carry their own completion through expect_tx + complete-tx; MMA commits must be wired to a barrier by hand. Forgetting this wiring is its own deadlock — the epilogue waits on a barrier that no one will ever arrive on.

Multi-stage pipelines — one barrier per stage, not per iteration

A pipelined K-loop runs the same TMA → MMA handoff hundreds of times. The naive reading of "one barrier per handoff" would allocate hundreds of barriers — one per iteration. That is wrong, and it would exhaust SMEM. The phase bit exists precisely so you don't have to.

Instead you allocate a small fixed number of SMEM buffers (the "stages" of a software pipeline), a matching fixed number of barriers, and a small set of phase registers. Each logical iteration maps to one physical stage:

stage = iteration mod num_stages

The stage index selects which barrier and which buffer to use; the per-stage phase register (toggled each time that stage is reused) distinguishes the current use from the previous one. A two-stage pipeline is therefore exactly 2 mbarriers + 2 phase registers — not two hundred. Stage 0 handles even iterations, stage 1 handles odd ones, and each stage's phase toggles every time the loop comes back around to it.

This is the structure lesson 09 builds the PIPE_DEPTH=2 double-buffer ring on, and the reason barriers also signal resource reuse: a stage's SMEM buffer cannot be overwritten by the next load until that stage's consumers are done with it — which the barrier's next round enforces.

The three canonical handoffs

Across every modern kernel, mbarriers coordinate exactly three kinds of producer→consumer handoff. Two of them are barrier waits; the first needs an extra ingredient.

Handoff 1 is the one beginners miss because it looks like a normal barrier problem but isn't. The TMA store and the tensor core operate in a different memory "proxy" than ordinary thread stores; an mbarrier orders arrivals, but it does not by itself flush a thread's SMEM writes into view of the async proxy. A fence.proxy_async is what publishes those writes before the engine reads them. Skip it and the engine reads pre-write bytes — another silent-corruption path, distinct from the phase bug.