B200 Gemm Optimization Ladder

E — Executable steps (Execution)

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

1. Establish the baseline contract
   • Fix the math definition, layout, dtype, reference, timing framework, and representative shape set.
2. Level 1: single-tile correct path
   • Synchronous copy, one MMA, TMEM readback, store.
   • Acceptance: element-wise correct on small matrices, with every address space explainable.
3. Level 2: K-loop
   • Correctly handle the initial accumulator, per-K-tile accumulation, and the K tail.
4. Level 3: spatial tiling
   • grid→M/N tile mapping, boundary masks, full matrix coverage.
5. Level 4: TMA
   • descriptor/swizzle, load barrier, store drain; compare against the synchronous version bitwise/within tolerance.
6. Level 5: multi-stage pipeline
   • prologue/steady/epilogue, stage/phase ledger; measure the actual overlap.
7. Level 6: persistent scheduler
   • A fixed set of resident CTAs processes multiple tiles; check the tail and tile order.
8. Level 7: warp specialization
   • producer/MMA/writeback roles, the four classes of handoff, warpgroup-scoped sync.
9. Level 8: 2-CTA cluster
   • cluster tile, DSMEM/multicast, cta_group::2, remote barrier.
10. Level 9: multi-consumer

• Multiple consumers partition the N/M/output regions, ensuring the same staged operand feeds more compute without write conflicts.

11. Per-level gating

• Correctness: random shapes, misalignment, K=1/multi-tile, NaN/Inf, error thresholds.
• Performance: compare against the previous level, the library baseline, and the roofline.
• Resource: register/SMEM/TMEM, active clusters, spill.

12. Stopping rule

• If you are already close to the practical roofline or added complexity no longer brings stable gains, stop upgrading and keep the simpler version.

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

• cuBLASLt/CUTLASS already meets the need and there is no fusion, special dtype/layout, or research purpose.
• There is no correctness reference or reliable timing framework.

Failure modes

• Jumping multiple levels at once, making the source of an error impossible to localize.
• Testing only one neat large shape, ignoring small shapes and tail tiles.
• Treating warp specialization as a guaranteed speedup, ignoring the occupancy of the added roles and the sync.

Limitations

• The book's route is based mainly on TIRx and specific example shapes; when porting to CUDA/CUTLASS/Triton, keep the principles rather than copying the APIs verbatim.

Related skills

• depends-on: b200-kernel-roofline-triage, b200-scope-layout-dispatch
• contrasts-with: none
• composes-with: b200-tma-pipeline-designer, b200-mbarrier-protocol-auditor, b200-tcgen05-mma-contract-builder, b200-cluster-persistent-scheduler, b200-warp-specialized-debugger

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