Part II · the five hardware primitives

Clusters, distributed shared memory & cluster launch control

mbarriers let the engines inside one CTA coordinate. But one SM is too small to be the whole story: its shared memory can't stage a big-enough tile, neighboring CTAs reload the same operands, and a fixed schedule leaves SMs idle at the end. The fix is to make CTAs cooperate across SMs and to hand out work dynamically. This is the last primitive before we build a real GEMM.

These two are genuinely separate ideas that happen to ride on the same hardware unit (the cluster). Keep them apart as you read:

Thread-block clusters — the new scheduling unit

Through the whole track so far, the CTA (thread block) was the largest thing the hardware co-schedules: its threads share one SM and one shared-memory allocation, and that is the boundary. A thread-block cluster (introduced on Hopper, sm_90; extended on Blackwell) widens that boundary. A cluster is a small group of CTAs that the hardware guarantees are co-resident — launched together, possibly on different SMs within the same GPC, and able to synchronize and read/write each other's shared memory.

The cluster becomes the unit of launch. You declare its dimensions at launch time:

// host: cluster dims travel as a launch attribute, not a kernel arg
cudaLaunchConfig_t cfg = {};
cfg.gridDim  = grid;            // grid is now measured in CLUSTERS
cfg.blockDim = block;
cudaLaunchAttribute attr[1];
attr[0].id = cudaLaunchAttributeClusterDimension;
attr[0].val.clusterDim = { /*x*/ 2, /*y*/ 1, /*z*/ 1 };   // a 2-CTA cluster
cfg.attrs = attr; cfg.numAttrs = 1;
cudaLaunchKernelEx(&cfg, kernel, /*args…*/);

Inside the kernel, the cluster gets its own barrier — the analogue of __syncthreads() but across all CTAs in the cluster:

// CUTLASS / CuTe surface
cluster.sync();                 // wait for every CTA in the cluster

// lowers to a split arrive/wait, like an mbarrier:
barrier.cluster.arrive;         // this CTA has reached the cluster barrier
barrier.cluster.wait;           // block until all CTAs in the cluster have arrived

The split arrive/wait shape is deliberate: it is the same producer/consumer rhythm as an mbarrier, lifted from "threads in a CTA" to "CTAs in a cluster." Everything else in this lesson builds on that one capability — CTAs in a cluster can see each other.

Distributed shared memory — naming a neighbor's SMEM

Co-residence alone is not useful unless a CTA can actually address a peer's shared memory. That is distributed shared memory (DSMEM): within a cluster, a thread can name a location in another CTA's SMEM and move bytes into it directly — no detour through global memory.

The mapping is one instruction. Given a local SMEM address and the rank (id) of a peer CTA in the cluster, it returns the address of that same SMEM slot as seen in the peer:

// PTX: map a shared-memory address into a peer CTA's window
mapa.shared::cluster.u32   %peer_addr, %local_addr, %peer_cta_rank;

// the course's DSL spells the same thing:
//     map_shared_rank(buffer, cta_id)  →  the address of `buffer` inside CTA `cta_id`

Combine that remote address with the TMA bulk-copy engine from lesson 04 and you can stage a tile from HBM straight into a peer's SMEM, raising a completion barrier in that peer when the bytes land. The payoff is the headline of bottleneck (1): a tile staged once can be made visible to several CTAs without a global round trip. SMEM stops being a per-SM island and becomes a small shared pool across the cluster — which is exactly what we need to build an output tile larger than one SM could hold.

Two CTAs, one MMA — cta_group::2

Lesson 03 showed that tcgen05.mma takes a cta_group::1 / cta_group::2 modifier and left the second case as a promissory note. Here it pays off. In a 2-CTA cooperative MMA, two CTAs in a cluster each contribute their SMEM operands to a single, larger tensor-core tile:

• A is split by M rows. CTA-0 holds rows 0–127, CTA-1 holds rows 128–255 — logically the stacked operand [A0 ; A1]. Neither CTA holds the whole A.
• B is shared across the pair via DSMEM. One copy of the B slice is staged and made visible to both CTAs through the peer-SMEM mapping.
• The pair produces one 256×256 output tile, with each CTA owning its half of the accumulator in its own TMEM.
• Only CTA-0's elected lane issues the MMA. The hardware then drives the math across both SMs. The even CTA (CTA-0) commits the pair's completion barrier — one signal for the whole cooperative op.

Why this is the move that matters: each CTA's B slice is now multiplied against the other CTA's A slice as well as its own. The same loaded B bytes feed twice as much math — B's per-FLOP load halves, so its arithmetic intensity roughly doubles. (A is split across the pair either way, so the total tile AI rises by a smaller factor; the win is on the operand we were forced to keep re-reading.) On a machine whose roofline ridge sits near 250 FLOP/byte (lesson 00), pushing AI up is pushing the kernel toward the compute roof — and lesson 10 stacks a second consumer on top to halve B's load again.

TMA multicast — one load, many CTAs

Cooperation handles a pair sharing B. Bottleneck (1) had a wider form: in a tiled GEMM, every CTA computing a tile in the same grid row needs the same A operand, and every CTA in a column needs the same B. Issuing one TMA load per CTA reloads identical bytes from HBM many times over.

TMA multicast collapses that. A single load issued by the TMA engine delivers the same GMEM tile to several CTAs at once, selected by a bitmask of cluster ranks:

// one bulk-tensor load, broadcast to the CTAs named in the mask
cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes \
        .multicast::cluster                          \
        [dst_smem], [tensor_map, {x, y}], [mbar], ctaMask;
//                                                  ^^^^^^^  e.g. 0b11 = CTAs 0 and 1
// CUTLASS names this copy atom: SM90_TMA_LOAD_MULTICAST

The redundant global traffic — N copies of the same row operand — becomes one. Completion is still signalled the lesson-06 way: an mbarrier with an expect_tx byte budget in each receiving CTA. Multicast doesn't change how a load completes; it changes how many times the bytes cross the HBM bus. That is the second bandwidth win, and it is why the GEMM tile scheduler (lesson 09) deliberately groups tiles that share operands.

The second wall: the static schedule's tail

Now the temporal problem. The whole GEMM ladder (lessons 08–10) rides on the persistent-kernel pattern: instead of launching one CTA per output tile, you launch a fixed pool of CTAs sized to the machine — roughly the SM count — and each one loops over many tiles, keeping operands hot in L2 and amortizing launch cost over the whole problem. The classic version computes "my next tile" from a formula:

// static persistent schedule (the thing CLC replaces)
for (int tile = blockIdx_in_pool; tile < num_tiles; tile += pool_size) {
    compute_tile(tile);            // next tile = a fixed stride. decided BEFORE any work runs.
}

This is correct and it is fast — until the tiles cost different amounts. In a causal-attention problem, or any GEMM where num_tiles isn't a clean multiple of the pool size, the per-SM tile counts come out uneven. The SMs that drew cheap or few tiles finish early and idle; wall-clock is set by the unluckiest SM. That idle stretch at the end is the tail, and a static schedule cannot fix it, because the assignment was frozen before the kernel knew how long anything would take.

Cluster launch control — hardware work-stealing

Cluster launch control (CLC) is the Blackwell mechanism that lets a resident cluster ask the hardware for another tile at runtime. Instead of computing "my next tile = formula," a cluster that is about to run dry issues a request to claim a not-yet-launched cluster's work. If the grid still has pending clusters, the request succeeds and hands back a coordinate; the early-finishing cluster picks up that work instead of idling. It is work-stealing, arbitrated in hardware.

Crucially, CLC speaks the exact same async protocol as TMA from lesson 06 — an async request whose completion is reported through an mbarrier, with a phase bit you track. Three instructions do the work:

// 1) fire an ASYNC request to cancel-and-claim a pending cluster's launch.
//    the 16-byte response lands in SMEM; completion is signalled on an mbarrier
//    (the SAME barrier-and-phase model as a TMA load — see lesson 06).
clusterlaunchcontrol.try_cancel.async.shared::cta.mbarrier::complete_tx::bytes \
        [response_smem], [mbar];

// 2) did the steal succeed? returns a predicate.
clusterlaunchcontrol.query_cancel.is_canceled.pred.b128   %got, response;

// 3) if it did, decode the stolen cluster's CTA id → a coordinate (x,y,z).
//    VALID ONLY when is_canceled is true.
clusterlaunchcontrol.query_cancel.get_first_ctaid.v4.b32.b128  {%x,%y,%z,%w}, response;

Note what the response is: a 16-byte record written to shared memory, and an mbarrier arrival that flips a phase exactly like a completed TMA load. You already know how to wait on this — it is lesson 06's machinery, reused. CLC didn't invent a new synchronization model; it plugged work-distribution into the one we already have.

The work-stealing loop — "ask first, then compute"

The loop has a specific shape, and the ordering is the whole point. You issue the request for the next tile before you compute the current one, so the scheduler's round-trip overlaps with useful math instead of stalling at the end:

1. Issue try_cancel for a possible next tile — fired off before the current tile's math begins, so the request is in flight while the tensor core works.
2. Compute the current tile while the request travels to the scheduler and back.
3. Wait on the response mbarrier — same try_wait + phase toggle as a TMA completion (lesson 06).
4. Query is_canceled — did we successfully steal a pending cluster's work?
5. Branch on the predicate. If true: decode get_first_ctaid → (x, y, z) → that becomes the next output tile; loop back to step 1. If false: there is no more work — exit.

Two things to nail down, because both are easy to get wrong:

• There is no sentinel tile id. The loop does not look for a magic "−1" or "out of range" coordinate. It ends purely on the is_canceled predicate coming back false. The coordinate from get_first_ctaid is meaningful only when the predicate was true — reading it otherwise is garbage.
• Ask-then-compute is the overlap. If you computed first and asked afterward, the scheduler latency would land in series and reopen a (smaller) tail. Issuing the request first hides that latency behind the current tile's math — the same "issue the async op, then do work, then wait" pattern TMA taught us.

The result: a cluster that finishes early doesn't wait for a static stride to grant it more work — it pulls the next pending cluster coordinate the instant it's free. The tail collapses toward the cost of a single tile, and all SMs finish close together.

Where these plug into the GEMM

Both halves of this lesson are bottlenecks we deliberately planted for the ladder ahead, and both get spent there:

• CLC rides the persistent kernel. Lesson 09 builds the persistent kernel and its tile scheduler with the static "next tile = formula" step; CLC is the drop-in upgrade that replaces that formula with "ask the hardware," killing the tail.
• The 2-CTA cluster is a GEMM step of its own. Lesson 10 turns the single-CTA warp-specialized kernel into a cta_group::2 pair with DSMEM + multicast — the 256×256-tile, AI-doubling step that pushes the kernel to the compute roof.