[lucifer1004]

Status

Stacked investigation draft. Depends on #1236 and tracks the possible second slice of #1235. The head currently contains only the dependency stack; no counted-HMMA implementation has been added yet.

The current evidence does not yet justify a carrier-formation pass. This draft remains open while the engine-side x4 schedule is exhausted.

Corrected problem statement

The original Rust kernel encoded one N8 BF16 B fragment as one ldmatrix.x2. That was the wrong abstraction boundary, not a compiler combine failure. Every admitted general-BF16 topology has an even number of N8 fragments per warp, so adjacent fragments can be loaded as x4 pairs with no x2 tail.

Matched NCU disassembly for M8192/N1024 had shown:

• cuBLAS nvjet_sm120_tss_mma_128x176x64_2_32x88x64_tmaAB_alignCD4_bz_TNNN: 30 static LDSM.16.M88.4, no x2, and 88 HMMA;
• old Oxide BM128/BN128/S2: 8 static LDSM.16.M88.4, 32 LDSM.16.M88.2, and 64 HMMA.

After pairing adjacent B fragments in ordinary Rust, all 20 registered general-BF16 variants build with x4-only SASS. Representative static bodies are:

• BM128/BN64: 16 x4, 0 x2;
• BM128/BN128: 24 x4, 0 x2;
• BM64/BN64: 12 x4, 0 x2;
• BM64/BN128: 20 x4, 0 x2;
• BM32/BN64: 8 x4, 0 x2.

All 20 variants passed the repository patterned-BF16 correctness/clock case. At M8192/N1024, the best BN128 median improved from the previous approximately 207.232 us to 196.608 us. At M8192/N64, the best BN64 median was 32.544 us versus the previous approximately 32.384 us, which is not a meaningful change.

Changing the x4-only body from pointer addresses to the #1236 .shared::cta.u32 API did not materially change timing. It preserved the no-spill resource shape for BM32/BM64 and BM128/BN64. BM128/BN128 remains on a register-allocation cliff: x4-only pointer lowered to REG168/STACK56; x4-only shared-u32 lowered to REG168/STACK48. Despite that spill, the x4 schedule is faster than the old x2-heavy control.

The earlier isolated two-x4/eight-x2-per-step spike proves only that a u32 address carrier can survive both compiler routes. It is not a production micro-schedule and is no longer an acceptance target.

Remaining question

Can the BM128/BN128 x4-only Rust schedule be reordered or represented more directly so it returns to zero spill without a new compiler carrier?

Only if source-level scheduling cannot remove that final spill while preserving the measured x4 performance should this PR introduce a compiler facility. Such a facility would preserve a kernel-authored counted register lifetime through libNVVM; it would not select x4/x2 grouping, tile geometry, HMMA issue order, or tuning policy.

Non-goals

• restoring or admitting the old x2-heavy schedule
• a public application-specific GEMM entry point
• automatic kernel fusion
• a general scheduling DSL or arbitrary PTX region
• tuning policy or hardware-specific candidate selection
• replacing the existing individual register-MMA and ldmatrix operations
• changing WGMMA counted-loop behavior

Evidence required before any implementation direction

• exhaust the small set of source-level x4 issue-order/live-range alternatives on BM128/BN128;
• retain x4-only SASS and repository correctness;
• compare REG/STACK/local load-store counts and matched repository medians;
• demonstrate that an ordinary Rust representation cannot meet the no-spill gate before adding a pass or compound carrier.

Draft discussion will revise or close this PR based on that evidence.

[nihalpasham]

Closing as a duplicate of #1236: both branches point at 4284c72, so nothing here is not already there, and the carrier pass it was opened for is not written. Slice-2 discussion goes on #1235.