B200 Gemm Optimization Ladder

B200 GEMM Optimization Ladder

R — Source evidence (Reading, paraphrased)

• [S11] Start from a minimal correct single tile, then add K accumulation and multi-CTA spatial tiling one step at a time, avoiding debugging all the complexity at once.
• [S12] First let TMA take over the regular tile copies, then use multi-stage SMEM and a persistent scheduler to reduce waiting.
• [S13] Warp specialization assigns load, MMA, and writeback to different roles; 2-CTA cluster and multi-consumer further remove serial bottlenecks.
• [S3] Every level should be driven by roofline and measured evidence; a more complex structure is not guaranteed to be faster.

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)

GEMM optimization is not generating the "final kernel" in one shot, but a regression-testable, bisectable upgrade path. Every level must simultaneously satisfy:

• correct against the reference;
• an explainable layout/synchronization contract;
• the performance change is measured;
• if it regresses, you know where the added resource or serialization point is.

The agent should keep the version at every level; jumping straight to a complex warp-specialized cluster kernel and then blindly guessing at bugs is forbidden.

A1 — Applications in the source (Past Application)

The nine-level route in the book

1. A single 128×128 output tile.
2. K-loop accumulation.
3. Multiple CTAs covering the full M/N.
4. TMA async load/store.
5. PIPE_DEPTH=2 software pipeline.
6. Persistent kernel + tile scheduler.
7. Warp specialization.
8. 2-CTA cluster.
9. Multi-consumer warp specialization.

Every level keeps the same basic data path: GMEM→SMEM→tcgen05→TMEM→register/SMEM→GMEM, changing only concurrency and scheduling.

A2 — Trigger scenarios (Future Trigger) ★

In what situations will the user need this skill?

1. "Start from a correct GEMM and optimize it step by step into a high-performance B200 version."
2. "My GEMM is at Step 5 now; should the next step be persistent or warp specialization?"
3. "Write correctness and performance acceptance criteria for each level."

Language signals

• "Start from a correct GEMM and optimize it step by step into a high-performance B200 version."
• "My GEMM is at Step 5 now; should the next step be persistent or warp specialization?"
• "Write correctness and performance acceptance criteria for each level."

Distinction from adjacent skills

Difference from b200-kernel-roofline-triage: this skill is the GEMM-specific implementation route; the roofline skill decides whether compute/overlap optimization should continue to be pursued. Combine with the individual specialized skills to complete the concrete stages.

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