Repository navigation
Add two-slot F16 WGMMA partial-wait pipeline lowering #1144
Description
Summary
Add compiler support for the canonical straight-line F16 WGMMA partial-wait pipeline using two accumulator slots and wgmma.wait_group<1>.
F16 WGMMA currently supports canonical linear full-drain lowering and counted K-loop lowering, but partial-wait pipelines remain unsupported.
This issue is intentionally limited to the smallest useful F16 pipeline shape.
Target shape
Support the canonical straight-line sequence:
wgmma.fence
mma.f16 acc0
commit_group
mma.f16 acc1
commit_group
wait_group<1>
...
wait_group<0>
with:
m64n64k16.f32.f16.f16- exactly two independent accumulator slots
wait_group<1>for partial waits- a mandatory final
wait_group<0> - register-resident accumulators
Scope
In scope:
- F16
m64n64k16 - straight-line partial-wait pipelines
- exactly two accumulator slots
wait_group<1>- safe round-robin accumulator-slot reuse
- value-threaded lowering to tied
f32registers - backend validation with
ptxas
Out of scope:
- F16
wait_group<N>forN > 1 - F16 pipelines with more than two slots
- F16 counted pipelines
- TF32 partial-wait pipelines
- mixed BF16/F16 pipelines
- dynamic pipeline depths
- arbitrary CFG pipeline recognition
- non-canonical pointer fallback
Acceptance criteria
- Canonical two-slot F16
wait_group<1>pipelines are recognized and lowered. - The resulting PTX uses
wgmma.mma_async.sync.aligned.m64n64k16.f32.f16.f16. - Both accumulator slots remain register-resident.
- The final drain uses
wait_group<0>. - F16
wait_group<2>remains unsupported. - Mixed BF16/F16 partial-wait pipelines remain rejected.
- TF32 partial-wait pipelines remain unsupported.
- Backend validation reports zero spill loads and zero spill stores.
Activity
nihalpasham commentedon Aug 24, 2026
This stays on hold while #413 on
arch/pointer-contractis completed.This is not a host-side migration dependency: both branches edit the
mir-lowerWGMMA interface/value-carrier paths and overlapping compiler tests, so that compiler contract needs to settle first.Before we take this back up, please add:
ptxaschecks confirming both slots remain register-resident and the final wait drains the pipeline.ptxasacceptance and zero spills prove feasibility, not pipeline correctness or user value under the current acceptance criteria.uurl commentedon Aug 31, 2026
Closing this for now.
The implementation direction was validated in #1145, but the feature currently lacks both an external consumer and the Hopper hardware evidence required for acceptance: numerical validation of accumulator-slot reuse, an adversarial ordering case, and performance data against the full-drain path.
I don't currently have access to Hopper hardware to produce that evidence, so keeping this issue open would not be actionable.
If a concrete consumer appears in the future together with a path to the required sm_90a validation, this can be revisited from current
main.