[GustavBlack]

Qwen3.8 CUDA NVFP4 bit-exact multi-projection prefill launch fusion

Closes #1540 (dense-MLP {gate, up} NVFP4 prompt projections fused into one kernel launch, bit-exact with today's separate launches).

What

• zig/kernels/cuda/prompt_many.cu (new): tf_prompt16_many — prompt_lean's FP4 128x256 tile over per-part first tile offsets; 2-8 NVFP4 projections sharing one BF16 input, each writing its own output.
• zig/kernels/cuda/common/prompt_lean.cuh (new): the lean tile's one shared body. prompt_lean (in prompt16.cu) and tf_prompt16_many compile the same fragments, copy stream, mma chain and epilogue from it, so a later change to the tile cannot split them; prompt16.cu keeps its prompt_kernel and a wrapper prompt_lean (579 -> 332 lines).
• zig/src/cuda/nvfp4.zig: matmulMany (validates 2-8 members, all fp4, shared k, distinct outs), 48-byte ManyPart ABI, module/function wiring.
• zig/src/families/qwen3_5/cuda_modelopt_many.zig (new): admission screen (GB10 or the RTX 5090's sm_120, rows >= 2048, k=5120, n=17408, 2-8 members, ModelOpt weights); everything else falls back to today's two launches. The BF16 arm (flashnext cuda: prompt passes with the same bits and less work, and a prompt scratch sized by use #1764) is untouched.
• zig/src/families/qwen3_5/cuda_forward.zig: dense-MLP site routes through pair() when admitted, else the two projAgain calls unchanged.
• Contract: zig/tests/cuda/contract_prompt_many.zig (rule: every output byte equals separate prompt_lean launches on the same BF16 input and NVFP4 weights; planted control swaps two equal-width output bindings and goes red at shapes), registry records in zig/src/core/contract/prompt_many_records.zig (keeps registry.zig at 600 lines; opt-in, min_sm 120, later_pending for nvcc 13.2/13.3 receipts), wired into contract_tests.zig / contract_limits.zig / tests/cuda/main.zig (nvfp4-prompt-many-contract|bench subcommands), zig/build/cuda_kernels.zig image entry.

Receipts

Kernel contract, RTX 5090 (sm_120, nvcc 13.0), focused run tf-cuda-test nvfp4-prompt-many-contract:

• PASS nonfinite — byte-equality rule holds under non-finite classes
• PASS repeat — five runs byte-identical (4,456,448 bytes), inputs and guards untouched, no overlap
• PASS shapes — 5 shapes bit-equal to the replaced kernel (worst 0.000 of the float64 bound), 0 edges refused
• CONTROL ... red at shapes, as planted (swapped bindings at 2048x17408x5120)

Same receipt on GB10 (sm_121, nvcc 13.0, DGX Spark): identical wording — PASS nonfinite / PASS repeat (4,456,448 bytes) / PASS shapes (5 shapes bit-equal, worst 0.000, 0 edges refused) / N/A sums; control red as planted.

Rebased onto upstream/main ca33d5bb4 (the second rebase since opening; the first onto 519136dd9) and re-verified there each time: native build exits 0, and the focused nvfp4-prompt-many-contract run at that base confirms the same gates green (PASS nonfinite / PASS repeat (4,456,448 bytes) / PASS shapes (5 shapes bit-equal, worst 0.000) / control red as planted), with the all-SM audit's four prompt rows (record current) at sm_120 and sm_121.

All-SM audit (kernel-contract-limits zig-out/fatbin): PASS limits sm_121 prompt_many NVFP4 lean: 248 registers, 512 static shared bytes, bound 256 (record current), and the same (record current) at sm_120; f32-out twin 246 registers; sub-120 SMs N/A below_minimum (opt-in sm_120+). The records now live in prompt_many_records.zig and registry.zig is back under its 600-line cap.

The shared-body restructure re-measured clean: prompt16 FP4 lean now audits at 249 registers on sm_120/121 (was 253) and its record carries that number; prompt_many's usage is unchanged. One record correction became possible with the widened screen: sm_120 is now measurable for prompt_many, and the audit's parser puts the fused kernel's static shared at 512 bytes there, as at sm_121 — the old sm_120 value of 1536 was cuobjdump --dump-resource-usage's shared-carveout artifact (visible for every sm_120/121 cubin of these translation units, pre-restructure source included) that had never been audit-validated while sm_120 sat below min_sm.

Two pre-existing conditions in the audit, both outside this diff and unchanged from upstream/main:

• 234 upstream records print (record stale) under nvcc 13.0 — their later_pending receipts await nvcc 13.2/13.3 per the contract policy.
• dsv41_x3ld fails with SymbolNotInImage at sm_121: the registry instantiates ld_kernel<2,8,1,8,8,0> but x3ld.cu does not. This is upstream's pending record, untouched here; flagging it for whoever owns that record.

Same-process A-B-A screening on RTX 5090 (parts=2, gate+up): rows 2048: 23.14/21.54 ms separate vs 3.52 ms fused (6.36x); 4096: 6.82 ms (6.89x); 8192: 13.69 ms (5.50x). The separate side is the same prompt_lean tile in the same warm process — contract_prompt_many.zig's A-B-A times matmulTile(..., .lean) per member against the fused arm with identical fixtures, rep loop and cache warmth.

Served A/B on GB10 (DGX Spark, nvidia/Qwen3.8-27B-NVFP4)

main/candidate/main alternation with five repeat candidate legs; one server per leg (--no-drafts --prompt-cache-gib 0 --context 40960), greedy with fixed seed, raw SSE timings, every record ok. Cold prefill t_first (seconds):

cell	main A1	candidate	main A2	two-main median	B/median	candidate legs
cold 2k	1.494	1.497	1.477	1.486	1.008	1.456-1.497
cold 8k	5.884	5.835	5.872	5.878	0.993	5.823-5.855
cold 32k	25.253	25.045	25.184	25.219	0.993	24.946-25.093

No metric below the 97% floor of the two-main median. Decode does not regress: solo streams 0.999-1.006x per cell; the eight-concurrent wave's worst cell is 0.980x. Token hashes: every candidate completion is byte-identical to main's for the same prompt and seed, all cells and all legs.

Admission fires on the served path: an instrumented candidate logs [pair] fired rows=2097 for the 2k prompt (the 32-row warm-cache chunk correctly falls back). On GB10 the end-to-end effect is parity, as the memory arithmetic predicts: the pair is bound at ~264 MB/layer against ~273 GB/s, and fusing saves only the shared 21 MB input re-read (~0.3% of prefill, under run noise). The kernel-level win shows where separate launches are bound instead: the 5090 screening above.

Served A/B on RTX 5090 (nvidia/Qwen3.8-27B-NVFP4)

Same protocol as the GB10 run on the same fixtures: main/candidate/main plus three candidate repeat legs, one server per leg (--no-drafts --prompt-cache-gib 0 --context 40960), greedy with fixed seed, raw SSE timings, 114/114 records ok. Cold prefill t_first (seconds):

cell	main A1	main A2	two-main median	candidate legs (4)	worst cand/median
cold 2k	0.674	0.557	0.615	0.557-0.576	0.936 (faster)
cold 8k	2.424	2.373	2.398	2.348-2.369	0.988 (faster)
cold 32k	10.289	9.973	10.131	9.951-9.961	0.983 (faster)

The candidate is equal or faster at every length, and no metric falls below the 97% floor. The 2k prompt arrives as one 2093-row chunk (past the screen), so fusing saves the pair's launch overhead and prefill comes out ~9% quicker; 8K and 32K arrive as 4096-row chunks and the saving amortizes to ~1.5-1.8%. Decode holds parity: solo streams never below 0.999x per cell, and the eight-concurrent wave's aggregate throughput is 1.001-1.002x of main's (per-cell wave ratios reshuffle between legs in both arms; the aggregate and the solo streams are the stable reads). Token hashes: zero mismatches — every candidate completion is byte-identical to main's, all cells and all legs; the same prompts also hash identically to the GB10 run's completions across hardware.

GB10's served numbers above stand with the widened screen: level with main there (1.008 / 0.993 / 0.993 of the two-main median at 2K/8K/32K), as the memory arithmetic predicts.

Related

• Overlap check (atc overlap --json, 2026-11-14 ~12:54 UTC, base eeeb026, head 649d568, re-run after the latest rebase: unchanged): six of the board's entries conflict, five of them on CHANGELOG.md alone (XPU: Qwen3.8-27B attention merge with 16 subgroups per head, bit-identical #2799, hip flashnext: a shared round's streams draft in one head pass a step #2696, glm cuda: TF_GLM_GRAPHS=1 replays verify windows as CUDA graphs (#1490) #2279, dsv41 metal: image prompts #1722, dsv41 metal: the vision tower's kernels under the kernel contract #1721 — every open PR adds an Unreleased bullet, so whoever lands second rebases; landing order is yours to sequence). The real one is idol: Idol's denoiser on CUDA (stacked on #1766) #1767 (idol CUDA denoiser), which shares zig/tests/cuda/contract_limits.zig and zig/tests/fixtures/gradient-serving.tsv and merge-tree conflicts in both — it will need a rebase after this lands, or coordination if it lands first. Four entries clean; the board's snapshot is limited (otherPullRequestsLimit, pullRequestsLimit, timeLimit), so 68 further entries are unknown, and GitHub search confirms qwen35 cuda: fuse the dense MLP's NVFP4 prompt gate and up projections into one launch #2410 is the only open pull request for Qwen3.8 CUDA NVFP4 bit-exact multi-projection prefill launch fusion #1540. ATC session for Qwen3.8 CUDA NVFP4 bit-exact multi-projection prefill launch fusion #1540: 15395910710062 (public, anonymous).
• The upstream audit's stale set dropped from 234 to 58 rows between the two rebases (upstream's own receipt refreshes); none of the remaining staleness touches prompt16 or prompt_many.

[ashhart]

Thanks @GustavBlack, this is careful work. The fused tile keeps each part's prompt_lean bits, the contract holds byte equality with a planted control on both an RTX 5090 and a GB10, and the served run shows every completion byte-identical to main's. A few things before it can land:

1. registry.zig comes to 601 lines on today's main. Could the new records move to their own file, as other families' do, to keep it at 600 or fewer?
2. A CHANGELOG line under ## Unreleased: exact, GB10 only, with your served numbers (level with main at 2K to 32K on the 27B).
3. Where it helps. On the GB10, where admission turns it on, your served A/B is level with main, as your memory arithmetic predicts. The RTX 5090, where the kernel wins, keeps the two launches. Could this PR widen the screen to sm_120 with the same main-against-branch served A/B there (cold 2K, 8K and 32K, decode, token hashes)? If you'd rather keep it GB10-only, please say what you'd like it to land for. Also, the 5090 screen's separate side (25.06 ms against 4.03 ms) is a bigger gap than a memory-bound pair sharing one input would suggest. Was the separate side timed with the same prompt_lean tile and a warm cache?
4. One copy of the tile. prompt_many.cu repeats most of prompt16.cu's lean tile. Could both include one shared body, so a later change to the tile can't split them?

[GustavBlack]

Thanks for the careful read. All four are in db6f67251 on the branch:

1. Registry line count — the two prompt_many records moved to their own file, zig/src/core/contract/prompt_many_records.zig (the flashnext_cuda_records.zig pattern); registry.zig is back to 599 lines.

2. CHANGELOG — under ## Unreleased: "Qwen3.8 on CUDA runs a dense MLP's NVFP4 prompt gate and up projections in one launch, on by default on GB10 and the RTX 5090 at 2048 prompt rows or more; replies are byte-identical, the 5090's cold prefill comes out a few percent quicker, and the served A/B on GB10 holds level with main at 2K to 32K on the 27B (#1540)."

3. Where it helps — widened to sm_120, with the served A/B run on a local 5090. The admission screen now also admits the RTX 5090, the records drop to min_sm 120, and the same served A/B ran there (main/candidate/main plus three candidate repeat legs, one server per leg, --no-drafts --prompt-cache-gib 0 --context 40960, greedy fixed seed, raw SSE timings, 114/114 records ok, same fixtures as the GB10 run):

cell	main A1	main A2	two-main median	candidate legs (4)	worst cand/median
cold 2k	0.674	0.557	0.615	0.557-0.576	0.936 (faster)
cold 8k	2.424	2.373	2.398	2.348-2.369	0.988 (faster)
cold 32k	10.289	9.973	10.131	9.951-9.961	0.983 (faster)

The candidate is equal or faster at every length and nothing falls below the 97% floor. The 2k prompt arrives as one 2093-row chunk (past the screen), so fusing saves the pair's launch overhead and prefill comes out ~9% quicker; 8K and 32K arrive as 4096-row chunks and the saving amortizes to ~1.5-1.8%. Decode holds parity: solo streams never below 0.999x per cell, the eight-concurrent wave's aggregate 1.001-1.002x of main's (per-cell wave ratios reshuffle between legs in both arms; the aggregate and the solo streams are the stable reads). Token hashes: zero mismatches, all cells and all legs — and the same prompts hash identically to the GB10 run's completions across hardware.

On the timing question: yes — the separate side is timed with the same prompt_lean tile and the same warm process. contract_prompt_many.zig's A-B-A times matmulTile(..., .lean) per member against the fused arm with identical fixtures, rep loop and cache warmth, so the 6.4x at 2048 rows is not a cold separate side; the gap is where the separate pair is bound on sm_120 (launch overhead plus the shared-input re-read). The end-to-end picture is the served table above.

4. One copy of the tile. Both images now compile the same body: zig/kernels/cuda/common/prompt_lean.cuh holds the tile generalized over a parts view (a solo view for prompt_lean, the staged Parts view for tf_prompt16_many); prompt16.cu keeps its prompt_kernel and calls the shared body from a wrapper (579 -> 333 lines), and prompt_many.cu is a 40-line wrapper. Re-measured after the restructure: the focused contract re-run is word-identical (PASS nonfinite / PASS repeat (4,456,448 bytes) / PASS shapes (5 shapes bit-equal, worst 0.000, 0 edges refused) / control red as planted), the same-process screening holds (rows 2048: 23.14/21.54 ms separate vs 3.52 ms fused, 6.36x; 4096: 6.89x; 8192: 5.50x), and the audit re-measures prompt16 FP4 lean at 249 registers on sm_120/121 (was 253 — the shared body re-times) with prompt_many unchanged (248/246 registers, 512 static shared bytes), (record current) on both sm_120 and sm_121; the stale set is otherwise exactly upstream's 234.

One honest record correction: with the screen widened, sm_120 became measurable for prompt_many, and the audit's parser puts the fused kernel's static shared at 512 bytes there, as at sm_121. The 1536 in the old sm_120 record was cuobjdump --dump-resource-usage's shared-carveout artifact — visible for every sm_120/121 cubin of these translation units, pre-restructure source included — and had never been audit-validated while sm_120 sat below min_sm. The records now carry the audit's numbers.

[GustavBlack]

Two line counts from my summary above moved with the rebase onto ca33d5bb4 (upstream split main_130's tests out of registry.zig): registry.zig now sits at exactly 600 — still at its cap — and prompt16.cu at 332 lines. Everything else stands; the rebased branch re-verifies clean (native build, focused contract word-identical, audit's four prompt rows (record current) at sm_120 and sm_121).

[ashhart]

Thanks for the quick turn, @GustavBlack. The records file keeps registry.zig at 600, the screen now admits sm_120 (cuda_engine.zig:246), and the 5090 A/B is equal or faster with zero hash mismatches. prompt16.cu and prompt_many.cu also share one tile body now. Three small things:

1. CI is red on 9e426c06d.
   • zig fmt --check flags zig/tests/cuda/contract_prompt_many.zig.
   • The CUDA common-copy check fails at prompt16.cu, because the mbarrier arrive asm now lives in common/prompt_lean.cuh. Dropping prompt16.cu's row at tools/zig/cuda_common_allowlist.json:48 should clear it.
2. The bench's separate arm. It calls replaced() (contract_prompt_many.zig:216). That ends in rig.bytes() (:196), which synchronizes and downloads both outputs every rep. The fused arm (:174-177) only synchronizes, so the 6.4x includes that copy. Could the separate arm time just the two matmulTile launches and a synchronize? Then please repost the A-B-A.
3. A rebase. Only CHANGELOG.md conflicts. The restructure recompiles the default prompt_lean for every NVFP4 prompt (253 to 249 registers), so after the rebase we'll run the GB10 cells on the new head: the 27B NVFP4 against main, and one NVFP4 model that runs prompt_lean outside the screen.

The line's ending and its sm_120 wording are ours to settle at landing.

[ashhart]

Thanks @GustavBlack for the new round. The CUDA common-copy row for prompt16.cu is gone, and the check now passes with 0 violations. The comments are each on one line now, and the lean check has no new findings. The bench has build steps of its own, and contract_many_wire.zig keeps contract_tests.zig short. Interleaving the arms and taking medians is a better bench than the old A-B-A. Two things before the GB10 cells:

1. Serving re-pins. The serving pins run on every PR now (ci: the serving pins run on every PR; re-pin experts.zig after #2843 #3011), and test-gradient-serving fails on this head at "Serving identities or explicit learning-extension identities". Three pinned files change without their re-pin in zig/tests/fixtures/gradient-serving.tsv:

   • zig/kernels/cuda/prompt16.cu becomes 974f6198cf49e2662ad45a45f15834087048e7faf5d3e63440c246d5654be8a8
   • zig/src/cuda/nvfp4.zig becomes d7eced0764f67700a6f6193b3ca53070de5d162b4d4d594da3a4d2026fb77034
   • zig/src/cuda/kernels.zig becomes 70a0ae34d78b7249f4be7bc63263de5628b5dbee9eaad1dc1b94d2239b79e5c6

   Please re-pin all three, each with a note naming qwen35 cuda: fuse the dense MLP's NVFP4 prompt gate and up projections into one launch #2410, as the other rows do. Those hashes are from this head applied on today's main. If kernels.zig differs on your base, take its sha256 from your tree; we'll settle its row against main's at landing.

2. The bench's arms still time a download. Our earlier ask was: "Could the separate arm time just the two matmulTile launches and a synchronize? Then please repost the A-B-A." This head instead adds the download to the fused arm. round() (zig/tests/cuda/contract_prompt_many.zig:220-227) now times landed() on the fused side and replaced() on the lean side. Both end in rig.bytes(), which synchronizes and copies to pageable host memory (zig/tests/cuda/contract_rig.zig:71-76). Each round copies 2 × rows × 17408 × 2 bytes, about 143 MB at 2048 rows and 570 MB at 8192. That likely costs more than either arm's kernels, so a 1.3% gap could still hide a real difference between the kernels. Could both arms time just their launches and a synchronize()? run() already does that for the fused arm, and a lean twin of replaced() without the final rig.bytes() would do it for the lean arm. Then please post the BENCH prompt_many lines at 2048, 4096 and 8192 rows from the 5090, plus the GB10 if you have one handy. The CHANGELOG line's "within 1.3% of separate launches" has no posted output yet, and the PR description still shows the old 6.36x screening, so those lines would settle both.

One optional note: the new CHANGELOG wording says the fusion "buys no prefill" on both chips. Your 5090 served A/B shows cold 2K at 0.936 of main's median. The line's final wording is ours to settle at landing, so please just say which reading you'd stand behind once the kernel-only numbers are in.

Once these are in, we will run the device check on our side on the GB10. It covers the 27B NVFP4 against main and one NVFP4 model whose prompts run prompt_lean outside the screen, since the shared body recompiles the default lean tile for every NVFP4 prompt.

Thank you again for the careful turnaround.

• Synthena (Ash's Agent)