[uurl]

Closes #1144

Summary

Adds register-resident lowering for the canonical straight-line F16 WGMMA partial wait pipeline with exactly two accumulator slots and wait_group<1>.

This extends the existing BF16 partial-wait pipeline machinery to the narrowly scoped F16 m64n64k16.f32.f16.f16 case while preserving the existing fail-closed behavior for unsupported F16 pipeline shapes.

What changed

• Added an F16 value-threaded pipeline carrier for m64n64k16.
• Extended straight-line pipeline recognition to homogeneous F16 sequences.
• Restricted F16 partial-wait pipelines to:
  • exactly two accumulator slots
  • wait_group<1>
  • a mandatory final wait_group<0>
• Generalized straight-line WGMMA pipeline PTX generation across BF16 and F16.
• Kept counted-pipeline lowering BF16-only.
• Added positive and negative lowering coverage.
• Extended the backend ptxas probe to validate the F16 two-slot pipeline.
• Updated WGMMA support documentation.

Supported F16 shape

wgmma.fence
mma.f16 acc0
commit_group
mma.f16 acc1
commit_group
wait_group<1>
...
wait_group<0>

The accumulator slots are value-threaded through tied f32 registers and reused only after the corresponding pending group is safe to retire.

Intentionally unsupported

This PR does not add support for:

• F16 wait_group<N> where N > 1
• F16 pipelines with more than two accumulator slots
• F16 counted pipelines
• TF32 partial-wait pipelines
• mixed BF16/F16 pipelines
• dynamic pipeline depths
• arbitrary CFG pipeline shapes
• non-canonical F16 pointer fallback

Validation

Validated with:

cargo fmt --all -- --check
git diff --check

cargo test -p dialect-nvvm --test ops_test
cargo test -p mir-lower --test lowering_test
cargo test -p cuda-oxide-codegen --test wgmma_pipeline_carrier_ptx -- --nocapture

cargo clippy \
  -p cuda-device \
  -p dialect-nvvm \
  -p mir-importer \
  -p mir-lower \
  -p cuda-oxide-codegen \
  --all-targets -- -D warnings

Results:

dialect-nvvm:        51 passed, 0 failed
mir-lower:          129 passed, 0 failed
backend PTX probe:    2 passed, 0 failed

ptxas validation for sm_90a:

F16:
  0 bytes stack frame
  0 bytes spill stores
  0 bytes spill loads
  90 registers

BF16:
  0 bytes stack frame
  0 bytes spill stores
  0 bytes spill loads
  90 registers

The negative coverage also verifies that F16 wait_group<2> remains unsupported and that mixed BF16/F16 partial-wait pipelines are rejected.

[nihalpasham]

Putting this PR on hold until the pointer-contract refactor is stable and the remaining WGMMA correctness evidence is available.

• The two-slot wait<1> conditions are narrowly coherent in static review, but assembler acceptance does not establish source-level numerical correctness under slot reuse.
• The branch conflicts with current main and overlaps lowering work changing in the refactor.
• Review needs a current rebase, a real Hopper numerical before/after artifact, and revalidation of pointer-slot identity.

This remains a verification hold, not an approval. The required evidence follows the correctness standard in the Discord announcement.

[nihalpasham]

Went through this in depth as part of a full review pass over the open queue. Closing it, and I want to be clear about why, because the lowering work itself is not the problem.

What's good

• the carrier op, the fail-closed gates (wait_group<2> rejected, mixed bf16/f16 rejected), and the parametrized PTX template all follow the merged bf16 pipeline path correctly
• ptxas on sm_90a accepts the fused region with 0 spills

Why it doesn't clear the bar

New compiler features need a concrete consumer and evidenced benefit. That's the standing acceptance rule, and this PR has neither:

• issue Add two-slot F16 WGMMA partial-wait pipeline lowering #1144 was filed one minute before this PR, by the same author. A companion issue you file for your own PR isn't a consumer asking for the feature
• the hold comment asked for three things back in August: a real kernel that needs two-slot f16 partial waits, a numerical GPU test covering accumulator slot reuse, and perf numbers against the existing full-drain path. None arrived

And the core risk is still open. Zero spills proves the PTX assembles, not that the pipeline is correct:

wgmma.mma_async {acc0}; commit;     what ptxas checks:  syntax, registers, spills
wgmma.mma_async {acc1}; commit;     what nobody checked: does acc0 read back the
wgmma.wait_group 1;   <- reuse acc0    right values when its group retires while
wgmma.wait_group 0;                    acc1's group is still in flight on real HW?

Every test in the PR builds dialect IR by hand; nothing drives the f16 intrinsics from Rust source through the recognizer, and nothing ran on a Hopper GPU.

It also aged out under main

• feat(mir): preserve Rust pointer provenance with a typed fact oracle #1186 rewrote validate_pipeline_events around canonical accumulator identity, the function this PR edits
• refactor: split large files into per-concern module directories #1190 split tests/lowering_test.rs, the file this PR adds ~115 test lines to
• wgmma_mma_kind() already existed on your base; the hand-rolled two-way kind detection in collect_pipeline_events duplicates it

If you want to bring this back

Please resubmit a fresh PR for review that comes with: the workload, the sm_90a numerical test (slot reuse + one adversarial ordering case), perf vs full-drain, an end-to-end #[kernel] test, and the recognizer changes routed through wgmma_mma_kind() on current main. The design direction was never the issue.