B200 Tma Pipeline Designer

E — Executable steps (Execution)

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

1. Assess TMA suitability
   • Is the tile regular, rectangular, and describable by fixed strides?
   • Is the copy volume large enough to amortize the descriptor and synchronization overhead?
   • If not, keep the vectorized thread copy.
2. Define the tensor-map descriptor
   • Global shape/stride, element size, tile shape, coordinates, boundary policy, swizzle.
   • Completion criterion: the source address and the destination SMEM layout can be uniquely derived from the logical tile coordinates.
3. Choose the issuer
   • Designate one lane/thread to issue; do not let every thread redundantly issue the TMA.
   • Write the issue scope and the CTA consumption scope separately.
4. Define the load protocol
   • Initialize a barrier for each stage.
   • arrive.expect_tx(bytes) sets the expected bytes and completes the issuing thread's arrival.
   • Consumers read SMEM only after waiting on the correct phase.
5. Define the store protocol
   • After commit_group, use wait_group to guarantee outgoing stores have drained.
   • Do not mistakenly substitute the load-side byte-count protocol for store drain.
6. Design the stage ring
   • Give PIPE_DEPTH, per-stage A/B/temporary SMEM, the ready/free barriers, and the initial phase values.
   • Compute total SMEM; check whether it pushes residency to an unacceptable level.
7. Write the three-phase timeline
   • Prologue: fill the initial batch of stages.
   • Steady state: overlap load k+1 / compute k / store k-1.
   • Epilogue: stop issuing new loads, finish the remaining compute/store, drain the groups.
8. Validate
   • Correctness: asymmetric small matrices, boundary tiles, different K_TILES.
   • Protocol: each stage's ready and free each have a unique producer and consumer.
   • Performance: confirm the copy instructions genuinely overlap with Tensor Core work, rather than code that looks asynchronous but is still issued serially by the same role.

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

• Highly sparse, irregular gather/scatter that a descriptor cannot express.
• A one-off, tiny copy where the TMA setup and synchronization cost may exceed the benefit.

Failure modes

• TMA swizzle inconsistent with the MMA layout.
• Consuming right after the load is issued, missing the barrier wait.
• Reusing SMEM before the store has drained.
• Increasing pipeline depth without accounting for SMEM resources and occupancy.

Limitations

• The best pipeline depth must be measured on real shapes, clock frequencies, and compiled output; it cannot be derived from a fixed rule.

Related skills

• depends-on: b200-scope-layout-dispatch, b200-layout-contract-auditor
• contrasts-with: none
• composes-with: b200-mbarrier-protocol-auditor, b200-gemm-optimization-ladder, b200-flash-attention4-planner

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