B200 Cluster Persistent Scheduler

B200 Cluster & Persistent Scheduler

R — Source evidence (Reading, paraphrased)

• [S2/S16] CTAs within a cluster can access each other's SMEM (DSMEM); the B200 portable cluster size is 8, and an explicit opt-in reaches the non-portable 16, but this may reduce the number of active blocks.
• [S13] A 2-CTA cluster can jointly compute a larger MMA tile and share operands through cluster-scope handoff.
• [S10] A persistent kernel keeps a fixed set of CTAs/clusters resident to process multiple tiles; CLC allows "stealing" work at runtime from not-yet-launched cluster coordinates, improving the tail.

Source: distilled from "Modern GPU Programming for MLSys" (https://mlc.ai/modern-gpu-programming-for-mlsys/) and the NVIDIA Blackwell tuning/compatibility guides. Short paraphrases only; no long passages are reproduced.

I — Methodology skeleton (Interpretation)

Cluster and persistent scheduling solve two different problems:

• cluster enlarges the spatial extent and data-sharing scope of a single cooperative tile;
• persistent scheduling reduces the static one-tile-one-CTA binding, letting a limited set of work owners fetch tasks in a loop;
• CLC on Blackwell further turns task fetching from static grid-stride into hardware-assisted dynamic tail scheduling.

Before using them, you must prove the gain comes from higher reuse, larger MMAs, or a better tail — not from "Blackwell has this feature, so use it".

A1 — Applications in the source (Past Application)

Case 1: 2-CTA GEMM

• The two CTAs each hold part of A/B/accumulator.
• DSMEM or multicast gives the pair the operands they both need.
• The barrier's CTA mask changes from single-CTA to cluster remote notification.

Case 2: uneven tile tail

• Static grid-stride may leave some SMs idle early during the final phase.
• CLC allows a finished resident cluster to cancel a not-yet-launched cluster and take over its coordinates to continue working.

A2 — Trigger scenarios (Future Trigger) ★

In what situations will the user need this skill?

1. "Is a 2-CTA cluster worth it for this GEMM?"
2. "Help me design a persistent tile scheduler and CLC tail stealing."
3. "How do I reason about DSMEM access, cluster size, and occupancy together?"

Language signals

• "Is a 2-CTA cluster worth it for this GEMM?"
• "Help me design a persistent tile scheduler and CLC tail stealing."
• "How do I reason about DSMEM access, cluster size, and occupancy together?"

Distinction from adjacent skills

Difference from b200-tcgen05-mma-contract-builder: this skill covers cluster-level cooperation and task scheduling; the MMA builder covers the contract of a single Tensor Core operation. Difference from b200-gemm-optimization-ladder: this skill also applies to non-GEMM persistent/cluster kernels.

E — Executable steps (Execution)

Once the skill is activated, the agent must execute the following process:

1. Prove the necessity of the cluster
   • Is a single CTA limited by SMEM/TMEM/operand reuse?
   • Does a 2-CTA tile raise Tensor Core utilization or reduce HBM/L2 traffic?
2. Define the cluster shape and tile mapping
   • clusterDim, each CTA's logical coordinates, the M/N/K slices, owners, and shared operands.
3. Plan DSMEM/multicast
   • Make accesses as coalesced as possible and 32B-aligned; avoid non-unit strides.
   • Write out the producers/consumers and fences for local SMEM and remote SMEM.
4. Compute residency
   • Use cudaOccupancyMaxActiveClusters or an equivalent tool.
   • Compare portable size ≤8 with the B200 opt-in 16; if using 16, set the nonportable attribute and flag the portability.
5. Upgrade the barrier scope
   • Update the CTA mask, remote arrival, and pair completion; check the even/leader CTA responsibilities.
6. Choose the scheduling mode
   • Equal-cost, uniform work: static tile formula/grid-stride.
   • Uneven cost or a pronounced tail: persistent owner + CLC work stealing.
7. Implement the CLC protocol
   • async try-cancel → mbarrier completion → query predicate → read the cluster coordinate only on success.
   • No numeric sentinel; termination is decided by the predicate.
8. Evaluate the tail and fairness
   • Test the boundary cases where the tile count is fewer than, equal to, and slightly more than the resident clusters.
9. Output the fallback
   • Provide the portable single-CTA/static-scheduler path and the conditions for enabling it.

Required outputs

1. Conclusion: the current choice/diagnosis; do not use a vague "it could be any of them".
2. Evidence or assumptions: which items come from user data, and which are hypotheses awaiting verification.
3. Contract/table/timeline: the auditable intermediate artifacts corresponding to this skill.
4. Minimal validation: correctness tests, boundary tests, and one falsifiable experiment.
5. Risks and fallback: alternative paths when hardware, version, or resource requirements are not met.

B — Boundaries (Boundary) ★

Do not use when

• A single CTA already saturates and tile costs are uniform; a cluster would only add synchronization and resource pressure.
• It must run fully identically on non-B200/non-Blackwell platforms, yet there is no fallback.

Failure modes

• Assuming a larger cluster size is always faster.
• DSMEM accesses that are not coalesced or not correctly fenced.
• Still reading an invalid coordinate when CLC fails.
• A persistent pool that is too large or too small, ignoring resource residency.

Limitations

• CLC and non-portable cluster sizes are architecture-dependent features; production code must have device-capability checks and a fallback.

Related skills

• depends-on: b200-scope-layout-dispatch, b200-mbarrier-protocol-auditor
• contrasts-with: none
• composes-with: b200-tcgen05-mma-contract-builder, b200-gemm-optimization-ladder, b200-kernel-roofline-triage

Audit info

• Validation passed: V1 ✓ / V2 ✓ / V3 ✓
• Test definitions: 6 (3 should_trigger / 2 should_not_trigger / 1 edge_case)
• Hardware validation: not performed; must be verified on a target B200
• Distilled: 2026-06-25