B200 Tcgen05 Mma Contract Builder

B200 tcgen05 MMA Contract Builder

R — Source evidence (Reading, paraphrased)

• [S7] tcgen05 is the Blackwell Tensor Core instruction family, issued by a single elected thread on behalf of the participating group; the operation itself is asynchronous.
• [S7] A/B usually reside in SMEM and the accumulator in TMEM; cta_group::2 makes the two CTAs of the same cluster cooperate, each keeping its own accumulator fragment.
• [S7] The block-scaled mode adds SFA/SFB: the data stays in SMEM, while the scale factors are supplied through TMEM.
• [S16/S17] B200 is compute capability 10.0; when using architecture-conditional features, make the portability boundary explicit.

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)

A complete MMA contract contains at least:

• the math shape M×N×K, the input/accumulate dtypes, and whether it is block-scaled;
• the participating unit, cta_group::1 or cta_group::2;
• which address space and layout each of A/B/scale/C lives in;
• who issues, who participates, and when it completes;
• the accumulator's mapping in TMEM and how the epilogue reads it back;
• boundary tiles, alignment, and the toolchain target.

If only a PTX mnemonic is given, without this contract, the agent should not generate code that merely "looks like it compiles".

A1 — Applications in the source (Past Application)

Case 1: cta_group::1, M=128

• A single CTA supplies the SMEM tiles of A/B.
• The 128 M rows map directly onto TMEM's 128 Lane rows; N maps onto Col.

Case 2: cta_group::2, M=256

• Two CTAs cooperate; each CTA owns 128 M rows and keeps the corresponding accumulator in its own TMEM.
• The even CTA is responsible for issuing the operation and for pair completion.

Case 3: block-scaled nvfp4/mxfp8

• The quantized A/B data is read from SMEM.
• SFA follows A's M partitioning; SFB, because both CTAs share B, must be visible/multicast to the pair.

A2 — Trigger scenarios (Future Trigger) ★

In what situations will the user need this skill?

1. "Choose the tcgen05 tile and cta_group for this B200 GEMM."
2. "Where should the scale factors of an nvfp4 block-scaled MMA go?"
3. "How is the accumulator split across the two CTAs' TMEM under cta_group::2?"

Language signals

• "Choose the tcgen05 tile and cta_group for this B200 GEMM."
• "Where should the scale factors of an nvfp4 block-scaled MMA go?"
• "How is the accumulator split across the two CTAs' TMEM under cta_group::2?"

Distinction from adjacent skills

Versus b200-tmem-lifecycle-planner: this skill first defines the MMA's math and hardware contract; the TMEM planner goes deeper into column budgeting and reuse. Versus b200-cluster-persistent-scheduler: this skill focuses on a single cooperative MMA; the latter focuses on cluster-level scheduling and tails.

E — Executable steps (Execution)

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

1. Confirm the target and toolchain
   • Target GPU, compute capability, whether sm_100a architecture-conditional features are allowed.
   • dtype and numerical-error requirements.
2. Define the math tile
   • M/N/K, transposes, accumulate semantics, boundary/remainder handling.
3. Choose the CTA group
   • group::1: simpler single-CTA resources and synchronization.
   • group::2: a larger cooperative tile and cross-CTA sharing, but added cluster/DSMEM/remote-barrier complexity.
4. Define operand placement
   • A/B SMEM layout, swizzle, and the slice each CTA holds.
   • For block-scaled, list the SFA/SFB shapes, the K block size, and the SMEM→TMEM copy.
5. Define the accumulator mapping
   • Write out each CTA's TLane/TCol formulas.
   • For modes such as M=64/128/256, make the lane packing explicit; avoid assuming a contiguous mapping.
6. Define issue and completion
   • The elected thread issues; the commit group is bound to a completion barrier.
   • Every TMEM consumer must run only after the barrier completes.
7. Define the epilogue
   • tcgen05.ld fragment shape, register cast/fusion, store path.
8. Run the three-contract check
   • The SMEM operand layout, the TMEM layout, and the async completion must all match.
9. Output the implementation skeleton and validation
   • Start with small shapes and random asymmetric data.
   • For dense and block-scaled separately, build a high-precision reference, error thresholds, and boundary K-block tests.

Required outputs

1. Conclusion: the current choice/diagnosis, without vague "could be any of them" hedging.
2. Evidence or assumptions: which 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 conditions are not met.

B — Boundaries (Boundary) ★

Do not use when

• The target is Ampere/Hopper, or only CUDA cores are used.
• The shape is too small and Tensor Core tile utilization extremely low; a custom MMA has not yet been shown to be worthwhile.

Failure modes

• Assuming every thread should issue the MMA.
• Treating the TMEM accumulator as a register fragment.
• Splitting only the compute under cta_group::2 without clarifying each CTA's operand/accumulator ownership.
• Wrong block-scale SFA/SFB layout or visibility.

Limitations

• The exact shapes/dtypes the instruction supports and the compiler APIs evolve; consult the current toolchain reference before generating code.

Related skills

• depends-on: b200-scope-layout-dispatch, b200-layout-contract-auditor
• contrasts-with: none
• composes-with: b200-tmem-lifecycle-planner, b200-mbarrier-protocol-auditor, b200-cluster-persistent-scheduler, b200-gemm-optimization-ladder

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