Repository navigation
nemotron cuda: ModelOpt checkpoints below sm_90, the lane matmul's slices reduced through a per-stream buffer - #901
BobClawblaw wants to merge 5 commits into
Conversation
…reduced through a per-stream buffer The lane matmul already reduces its K slices through fp32 partials where a card has no clusters (qmmf.plan's reduce, qmmf.cu's reduce_kernel), and its FP4 and FP8 decode to bf16 for mma.sync, but the family passed it no partials buffer and refused sm_89 and below at load. Each stream's scratch now holds the most partials any dense shape writes at rows short of fusing, Ops carries it to decode and prompt rows, and the refusal goes. The reduce adds the slices in the cluster path's order, so the sums do not change.
|
Thank you, @BobClawblaw, for adding the reduction buffer so ModelOpt checkpoints can run below sm_90. Two paths need attention. At Please wire each segment's own buffer and make the loader's Ops valid before scratch exists, then check short tails and the 255/256-row boundary on Ampere or Ada, including MTP, drafts on/off, resumed/fresh and concurrent/solo replies against the unchanged controls. We'd be happy to fold these fixes in ourselves at landing if you prefer. Posted by Ash's AI Assistant Synthena |
…and the loader's Ops needs no scratch A staggered prompt's segments built their Ops without a partials buffer, so below sm_90 a segment's rows short of fusing wrote the lane matmul's slices through a null pointer; each now passes its own scratch set's lane_part. The loader took Engine.ops(), which reads the buffers made after it; it now gets an Ops without one, since it uploads and packs and runs no lane matmul.
… without clusters reduces through The fp8 and bf16 lane checks ([320, 2048], split-K 4) passed no partials buffer, so at 16 and 128 rows below sm_90 the lane's reduce wrote through a null pointer; they now pass one sized for the widest unfused call, as the fp8 and nvfp4 tests do.
|
@ashhart thanks, both confirmed. Fixed at
The side stream's ChecksSetup. RTX 5090 (sm_120), Linux, CUDA 13.3,
The crash, before and after.
Segments. Ten prompts of 255, 256, 2,049, 2,303, 2,304, 4,097, 4,351, 4,352, 6,399 and 8,193 tokens, so one-row tails and both sides of the 255/256 boundary at 1 to 4 segments.
Served (
Tests.
Not run. A real Ampere or Ada card. The reduce build follows the same plan and kernels as one, but not sm_89's SASS. An update on the sm_89 receipt follows separately. |
|
@ashhart The sm_89 update: the same checks on a real Ada card. RTX 4090 (sm_89, no clusters), Windows 11, CUDA 13.3, The crash, before and after.
Boundaries. Templated chat prompts of exactly 254, 255, 256, 257, 511, 512, 2,047, 2,048, 2,049, 2,050 and 4,097 tokens, 48 greedy tokens each. Six server configurations: one segment, two segments, and drafts off, each with the prompt cache on and off.
Earlier on this card, at Tests.
One overlap with #836. #836 already gives those Not run: Ampere (sm_80/86), and Linux on Ada. |
1.0.12 names ModelOpt's floor (modeloptRefusal) at sm_90; this branch reduces split-K sums without clusters, so the floor is now sm_80, where FP8 and NVFP4 still decode to bf16 for mma.sync. lane_part follows main's six new scratch buffers, and covers the latent MoE's two projections, whose lane calls now pass it.
|
Merged main (1.0.12) at
Checks on the 5090, the same two builds as before (clusters off, and the branch as pushed):
The latent path was not run: no latent checkpoint reaches One note outside this PR: GLM's CUDA family ( |
…s its buffers' ck8 1.0.12's Forward.init rebuilds its Ops from the bound buffers to carry their e4m3 rows (ck8), which dropped the partials buffer the engine and each segment pass in: on a card without clusters the first prompt's lane matmul reduced through a null pointer (CUDA_ERROR_LAUNCH_FAILED on an RTX 4090). It now takes the same buffers' lane_part. tf-cuda-test modelopt runs its lane rows from sm_80, the floor this branch sets, instead of sm_90.
|
@ashhart One more fix after the 1.0.12 merge, found on the real sm_89 card: What broke. 1.0.12's
Checks at
Not run: the latent MoE path (no latent checkpoint fits this card) and Ampere. |
|
A correction to my merge note above. Its 5090 checks after the 1.0.12 merge did not run the reduce path. 1.0.12 reformatted Rerun on the 5090 with clusters really off (the edit checked before building):
|
|
Thanks @BobClawblaw for correcting the ineffective clusters-off edit and rerunning both sides. The actual failure before 0b60b39 and the corrected 56-reply comparison are the useful controls; the earlier merged-head run cannot support that reduction claim. We retained your segment, Forward and latent partial-buffer wiring, and added a common pre-launch refusal when a reduction needs missing storage. Host controls cover that guard without changing valid cluster/fused dispatch. Ampere and the latent checkpoint path remain unqualified by these receipts. The final integrated Windows and device runs still need their own results. Posted by Ash's AI Assistant Synthena |
2 participants
What this changes
Nemotron's ModelOpt checkpoints (NVFP4, FP8, bf16 lane rows) now load on cards below sm_90. Before this, they were refused at load with
ModeloptNeedsSm90. One commit onmain.The kernels already had what these cards need. The lane matmul decodes FP4 and FP8 to bf16 for
mma.sync, and its cluster code is compiled out below sm_90 (qmmf.cu). Without clusters,qmmf.plansums the K slices through fp32 partials andreduce_kernel. The family never passed it a partials buffer, so the load refused those cards instead.cuda_state.zig: each stream's scratch getslane_part. It holds the most partials any dense shape writes at rows short of fusing (qmmf.fused_rows - 1), computed with the lane's ownplan. That is 22 MB a stream for Nemotron 3.5 Lightning, set byo_projat split-K 8.cuda_kernels.zig:Opscarries the buffer to decode rows (dense) and prompt rows (prefillDense, for the bf16 and FP8G rows the lane takes). The sm_90 refusal goes.cuda_engine.zig:Engine.ops()passes the bound buffers'lane_part.The bits.
reduce_kerneladds the slices in the cluster path's order, slice 0 and then 1 to SK − 1, before the tensor scale. So a card without clusters makes the same sums as one with them. On sm_90 and later nothing changes except a buffer each stream doesn't use.Receipt
-Dsm=89.maindoes not build on Windows yet, so the receipts ran this commit on Windows: the native CUDA server, and Qwen3.5-2B and Qwen3.8-27B on CUDA #836's branch (5e1c3fea); the Nemotron sources it touches are the same asmain's. Checkpoint:nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4, with its MTP head drafting.the native CUDA engine cannot load … (ModeloptNeedsSm90).--context 16384 --parallel 2. Startup: model 19.51 GiB, two streams of 0.18 GiB, a 1.4 GiB prompt cache.bench_concurrent --levels 1,2 --alone --serial: 24 of 24 alone and 12 of 12 serial equal."draft": falseon five prompts, greedy and seeded T 0.7: 10 of 10 equal.bench_openai, 128 tokens): 468 code / 363 chat tok/s sampled, 553 / 414 greedy. Drafts accepted 3,977 of 6,153.bench_cache --require-cold, a cold 2,570-token prompt in 0.176 s (qualified).zig build test: 743 of 757 passed, 14 skipped, 0 failed.test-cuda-host: 469 of 474 passed, 5 skipped.tf-cuda-testquant-ref,modelopt,glueandsmokepass.quant-ref's NVFP4[320, 2560]takes split-K 4, so on this card it runs the reduce path: within float64 tolerance, and every row count's bits are equal.zig fmt --checkis clean. The lean check reports 528 problems, the same asmain.Another checkpoint it opens. #897's Nemotron 3.5 Bolt 74B-A9B stores its Mamba
in_projandout_projin NVFP4. With this change it serves on the 4090, text only, one stream:tools/zig/bolt_reference.py's fp32 forward.Checklist
"draft": false, a resumed prompt equals a fresh one, and concurrent equals solo.quant-ref's coverage on these cards, and before this change the load refused them.