Skip to content

fix(qwen3_5): re-verify 27B NVFP4 on the nvidia ModelOpt checkpoint — binding 0.72x, attributed (#164 follow-up) - #178

Merged
mudler merged 1 commit into
mainfrom
row/MODEL-QWEN35-GDN-FP8-INPROJ
Aug 9, 2026
Merged

fix(qwen3_5): re-verify 27B NVFP4 on the nvidia ModelOpt checkpoint — binding 0.72x, attributed (#164 follow-up)#178
mudler merged 1 commit into
mainfrom
row/MODEL-QWEN35-GDN-FP8-INPROJ

Conversation

@localai-bot

Copy link
Copy Markdown
Collaborator

Follow-up to #169. That PR fixed the reported symptom (BF16-only lm_head); running it against the real checkpoints on the GB10 box then uncovered two more layers behind it. Same root cause each time: the dense loader assumed BF16 wherever the compressed-tensors NVFP4 probe did not match.

Layer 2 — FP8 tower tensors

Both publishers quantize parts of the tower to FP8 and leave the rest BF16, and they disagree on which parts and on the scale layout:

checkpoint FP8 tensors scale
nvidia/Qwen3.6-27B-NVFP4 linear_attn in_proj_qkv / in_proj_z / out_proj per-tensor F32
unsloth/... @ccdaab7e whole tower per-output-channel BF16

One MaterializeBf16Source now sits under LoadBf16Direct, LoadBf16Transposed and LoadMergedBf16RawNK. BF16 stays zero-copy on the mmap (unchanged); FP8 is dequantized. Reading a per-channel scale as per-tensor would be silently wrong rather than loud, so the element count decides and anything else is rejected.

Layer 3 — ModelOpt NVFP4 naming

Our probe was has(<proj>.weight_packed) — compressed-tensors only. NVIDIA ships ModelOpt: <proj>.weight U8 + .weight_scale F8 + .weight_scale_2 F32. Every probe missed, so a genuinely NVFP4 tower fell through to the BF16 path and died at the first U8 tensor. IsNvfp4Projection now accepts both spellings; LoadNvfp4AnyNaming reads either, converting ModelOpt's direct scale to the CT divisor convention the shared dequant expects — the same conversion #169 made for lm_head.

W4A16 is the default for ModelOpt

These checkpoints carry a per-tensor input_scale. Consuming it sets alpha, flips IsTrueW4A4() and routes to the fp4-activation GEMM — which produced incoherent output on nvidia's checkpoint. Leaving alpha at 0 takes the weight-only dispatcher and generates correctly. VT_MODELOPT_W4A4=1 restores the other arm for A/B (documented in docs/ENVIRONMENT.md).

Verified on real weights, dgx GB10 (sm_121)

Against the real 21 GB nvidia/Qwen3.6-27B-NVFP4:

  • before: dense loader: expected BF16 for ...linear_attn.in_proj_qkv.weight
  • after: greedy "The capital of France is"" Paris."

Clean CUDA build on the box; clean CPU build; full scripts/check-*.py battery green.

This unblocks benchmarking the new NVFP4 publishes, which is the next step.

@mudler
mudler force-pushed the row/MODEL-QWEN35-GDN-FP8-INPROJ branch 2 times, most recently from fc549e4 to ba43904 Compare August 9, 2026 15:17
@localai-bot localai-bot changed the title fix(qwen3_5): load the 2026-08 Qwen3.6-27B NVFP4 republishes (#164 follow-up) fix(qwen3_5): re-verify 27B NVFP4 on the nvidia ModelOpt checkpoint — binding 0.72x, attributed (#164 follow-up) Aug 9, 2026
@localai-bot

Copy link
Copy Markdown
Collaborator Author

Rebased onto current main (2b08dd24) and rewritten. This now carries the whole re-verification, not just the loader fix.

Why it grew. unsloth/Qwen3.6-27B-NVFP4 re-quantized in place (@890bdef7 NVFP4, @ccdaab7e FP8 W8A8 under the same name), so every recorded 27B-NVFP4 number was tied to a revision nobody had pinned. nvidia/Qwen3.6-27B-NVFP4 @0893e160 is the publisher's own single-revision NVFP4 build and is the reference from here on.

Binding grid (warm servers, one flock, greedy, ignore_eos, 128 output tokens, medians of 3, vLLM 0.25.0 graphed):

c ours vLLM 0.25.0 ratio
1 8.76 12.29 0.713x
2 17.07 23.49 0.727x
4 33.01 45.42 0.727x
8 62.13 86.01 0.722x

We are behind, flatly. Greedy continuations are identical between engines. SACRED gates on the production build: test_qwen27_paged_engine 235/235, test_qwen36_paged_engine 315/315.

Two method defects are fixed here, both of which would have recurred:

  1. The previous sweep's build omitted -DVLLM_CPP_TRITON=ON, which #ifdefs out the entire Triton-AOT GDN family on a GDN-hybrid model. The SACRED gate refuses such a build; a hand-written bench script has no guard. Re-measured on a production build it is a wash, so the conclusion survived, but nothing off that build was quotable.
  2. Five checkpoint-gated tests, test_qwen27_paged_engine among them, resolved the snapshot by taking whatever fs::directory_iterator yielded first under a repo with two cached revisions. The filesystem was choosing which model the SACRED gate measured, and it has been green by luck. Now pinned via tests/parity/hf_snapshot.h.

Attribution. nsys defaults to --cuda-graph-trace=graph, which reports a captured decode graph as one range; the first profile read 0.030 ms/step against a 114 ms/token wall clock and made the sampler look like the cost of decode. At node granularity the step is 114.298 ms of GPU kernel time against 114 ms wall, so decode is ~100% GPU-busy and no host-side lever exists at c1. Breakdown: 40% NVFP4 MLP marlin (192 calls, ~68% of the bandwidth roof), 24.6% cuBLAS bf16 gemvx (112 calls, 28.2 ms for at most ~1.7 GB), 22.4% cuBLASLt FP8 nvjet, 10% a single lm_head GEMM.

The dense-marlin routing fork is closed here (and the hang that blocked it before is named: mirroring the shared call sequence drops a per-call workspace memset that the shared helper compensates for by zeroing at allocation, leaving marlin's spin-locks uninitialized). It is worth +0.5% at c2-c4, +0.8% at c8 — real, small, and it retires the prime suspect. The next lever is the bf16 gemvx line.

Known pre-existing preflight failure, not from this branch: commit-trailers on 39da53c7, a commit already merged to main by another author.

@localai-bot

Copy link
Copy Markdown
Collaborator Author

Gap partially closed: 0.72x → 0.85x. The gemvx line the profile pointed at turned out to be a loader bug, not a kernel one.

LoadAttnDense had exactly two branches — NVFP4 or LoadBf16RawNK — and so did LoadGdnDense's out_proj. A modelopt_mixed checkpoint quantizes the attention tower to FP8 W8A8 while leaving the MLP NVFP4, so those projections matched neither branch's intent and fell through to the bf16 path, which dequantizes: 1.562 GiB of self_attn + 1.41 GiB of linear_attn.out_proj became ~5.9 GiB of BF16, re-read every decode step and executed as cuBLAS gemvx.

This was never a missing capability. q/k/v/o_proj_fp8 and out_proj_fp8 already exist, MatmulFp8Cutlass* already consumes them, and LoadAttn on the MoE path has done exactly this since the 35B work — its own comment calls it the DEFAULT. Only the dense path never got the branch. vLLM keeps the same weights fp8 (modelopt.py process_weights_after_loading), so this is POL-MIRROR-VLLM.

c before after gain vs vLLM before → after
1 8.78 10.41 +18.6% 0.713x → 0.847x
2 16.83 20.22 +20.1% 0.727x → 0.861x
4 32.98 38.76 +17.5% 0.727x → 0.853x
8 62.29 72.49 +16.4% 0.722x → 0.843x

Peak host RSS 24.2 → 21.0 GiB. Per-token 114 ms → ~96 ms, ~78% of the bandwidth roof against vLLM's ~95%. The 32-token greedy continuation is byte-identical to vLLM's, captured from the same warm process as the numbers. SACRED gates on this build: 235/235 and 315/315.

Still behind. The remaining lever is the NVFP4 MLP marlin: 40% of the step, 192 calls, ~68% of roof.

@localai-bot

Copy link
Copy Markdown
Collaborator Author

35B: the "correctness defect" was a bit-exact oracle tie, and a router lever now puts decode at 0.98x.

Correctness — resolved, not fixed, because there was nothing to fix. One logprobs request settled what five configuration A/B arms could not. At the divergence position vLLM's own top-2 are the same float32 value:

pos top1 logprob top2 logprob diff
6 ' its' −0.011696805 ' iconic' −5.511696815 5.500
7 ' rich' (8807) −1.2221027612686157 ' iconic' (25438) −1.2221027612686157 0.0
8 ' history' −0.132903934 ' cultural' −2.507904053 2.375

vLLM's torch.argmax breaks the tie by lowest index. Our sampler already implements the same rule (ArgReduce, lowest-index-wins on the true global index), but our logits aren't bit-identical to vLLM's, so the equality doesn't reproduce on our side. Two earlier framings of this on the branch are superseded and recorded as such.

Speed — warp-shuffle MoE router top-k. The attribution put MoeRouterTopKKernel at 4.7% of the decode step: 19.5 µs for a few hundred comparisons, because the grid is one block per token and each of the k rounds did a block-wide tree reduction (~64 __syncthreads for k=8 over 256 experts). Barrier latency was the kernel.

c before after gain vs vLLM
1 70.58 73.24 +3.8% 0.945x → 0.980x
2 107.64 113.02 +5.0% 0.826x → 0.867x
4 194.87 199.90 +2.6% 0.952x → 0.977x
8 318.56 325.36 +2.1% 0.900x → 0.919x

Byte-identical by construction: top-k is an argmax reduction, and argmax under a total order is associative and commutative, so any reduction order gives the same (value, index). The softmax max/sum reductions in the same kernel are arithmetic and were deliberately left on their existing tree — reassociating those would move the denominator by an ulp.

Gates: test_ops_moe 33451/33451, test_ops_moe_grouped 440/440, test_ops_moe_grouped_bf16 19/19, test_qwen36_paged_engine 315/315, test_qwen27_paged_engine 235/235. Continuation unchanged from baseline.

c2 and c8 remain the weak cells and both carry the wider spread, so the honest next step there is tightening the measurement before chasing it. CastF32Kernel (3.1%) is the next named lever.

@mudler
mudler force-pushed the row/MODEL-QWEN35-GDN-FP8-INPROJ branch from e81f81e to c666d04 Compare August 9, 2026 18:31
…the gap

The 27B-NVFP4 grid of record was measured on `unsloth/Qwen3.6-27B-NVFP4`, and
that repo re-quantized in place: @890bdef7 is genuine NVFP4, @ccdaab7e is the
same name turned into FP8 W8A8. `nvidia/Qwen3.6-27B-NVFP4` @0893e160 is the
publisher's own single-revision NVFP4 build and is the reference from here on.
Re-measuring against it found us far behind, and most of that is now closed.

LOADING. Three defects, all the same shape: the dense loader assumed BF16
wherever the compressed-tensors NVFP4 probe did not match. One
`MaterializeBf16Source` now sits under the BF16 loaders (a per-channel scale read
as per-tensor is REJECTED rather than silently wrong); `IsNvfp4Projection`
accepts ModelOpt naming (`weight`/`weight_scale_2`) alongside compressed-tensors;
and W4A16 stays the default because consuming ModelOpt's `input_scale` flips
`IsTrueW4A4()` into the fp4-activation GEMM, which produced incoherent text here
(`VT_MODELOPT_W4A4=1` opts in).

THE GAP WAS ALSO A LOADER BUG. `LoadAttnDense` and `LoadGdnDense`'s `out_proj`
had only NVFP4-or-BF16 branches, so a `modelopt_mixed` FP8 tower fell through to
the BF16 path, which DEQUANTIZES: ~3 GiB of FP8 became ~5.9 GiB of BF16 re-read
every decode step and executed as cuBLAS `gemvx`. Never a missing capability --
the `*_fp8` slots and their `MatmulFp8Cutlass*` consumers already existed and the
MoE loader has done exactly this since the 35B work, with its own comment calling
it the DEFAULT. vLLM keeps the same weights fp8, so this is POL-MIRROR-VLLM.

  27B, warm servers, greedy, ignore_eos, 128 tokens, medians of 3, vs vLLM 0.25.0
  c  | before | after | vs vLLM before -> after
  c1 |   8.76 | 10.41 | 0.713x -> 0.847x
  c8 |  62.13 | 72.49 | 0.722x -> 0.843x
  peak host RSS 24.2 -> 21.0 GiB

TWO ROUTER KERNELS WERE BARRIER-BOUND, not compute-bound. The router grid is one
block per token, so at decode a single block ran the k selection rounds with a
block-wide tree each (~64 `__syncthreads` for k=8 over 256 experts); the GROUPED
variant was worse, running group scoring, mask, top-k and renorm on ONE lane
behind `if (threadIdx.x != 0) return;`, which its own comment deferred as "W9".
Both now use a warp-shuffle argmax. Byte-identical by construction: top-k is an
ARGMAX over a total order, which reassociates freely, while the softmax max and
sum reductions in the same kernels are ARITHMETIC and are deliberately left on
their original trees.

  35B-A3B @491c2f1e, same harness: c1 70.58 -> 73.24 (0.945x -> 0.980x),
  c4 194.87 -> 199.90 (0.952x -> 0.977x), c2 0.826x -> 0.867x, c8 0.900x -> 0.919x

THE 27B GATE WAS LETTING THE FILESYSTEM PICK THE MODEL. Five checkpoint-gated
tests, SACRED `test_qwen27_paged_engine` among them, took the first entry
`fs::directory_iterator` yielded under a repo with TWO revisions, so a token-exact
pass against the FP8 revision would have been recorded as an NVFP4 pass. Green by
luck. `tests/parity/hf_snapshot.h` pins the revision the goldens' own
`oracle.model` field names, with `VT_QWEN27_SNAPSHOT` as the escape hatch.

`max_num_seqs` also moves 8 -> 32: at 8 a c8 client sat exactly on our own batch
ceiling. Not vLLM's 1024, because it also caps the padded decode-graph set.

Correctness: 27B greedy continuation byte-identical to vLLM across 5 repeats.
The 35B's continuation differs at ONE token, and that is NOT a defect -- vLLM's
own top-2 there are the same float32 value (-1.2221027612686157, diff 0.0) and
its `torch.argmax` breaks the tie by lower index; our `ArgReduce` implements the
same rule, but our logits are not bit-identical, so the tie does not reproduce.

Gates: test_qwen27_paged_engine 235/235, test_qwen36_paged_engine 315/315,
test_ops_moe 33451/33451, test_ops_moe_grouped 440/440, grouped_bf16 19/19,
test_deepseek_v4_moe 716/716.

Still behind, and named rather than hidden: 27B at 0.85x with the NVFP4 MLP
marlin at ~68% of the bandwidth roof, 35B c2/c8 the weak cells, `CastF32Kernel`
3.1% of the 35B step. Full method, the refuted hypotheses and the decode
attributions are in `.agents/benchmark-record.md`.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-opus-5 [ClaudeCode]
@mudler
mudler force-pushed the row/MODEL-QWEN35-GDN-FP8-INPROJ branch from c666d04 to 6603356 Compare August 9, 2026 19:51
@mudler
mudler merged commit 0cf3dbb into main Aug 9, 2026
9 of 13 checks passed
mudler added a commit that referenced this pull request Aug 9, 2026
…message

Main CI has been red on EVERY merged PR. Two independent defects, both landed
by #178's push (run 31332846716); `cuda-fat-build` in that run was cancelled by
concurrency, not failing.

THE GATE REDDENED MAIN FOR OBEYING IT. `check-role-discipline.py` judges every
commit in the `before..after` range on its own message. A PR landed with a REAL
merge commit pushes the merge AND the branch commits under it: the merge names
the PR, the branch commits were never required to, so each one read as a direct
push. `6603356a` (#178), `e73cbbae` (#204) and `1a02ab4f` (#196) all failed this
way, in both `documentation-checkpoint` and `agent-record`. Arrival is now judged
ONCE, on the commit that lands the change: `merged_pr_content` exempts what a
row/* PR merge brings in. Squash-merges are untouched -- their one commit carries
"(#N)" and passes on its own message. NOT a weakening, and gated as such: only
the SIDE parents count, so `--not parents[0]` keeps a commit pushed straight to
main from being laundered by merging a PR on top, and a merge naming no row and
no PR exempts nothing. Four unit checks build real git history for those cases,
plus the exact `3bbee96e..0cf3dbb` range CI ran, pinned with `has_reached_main`
forced TRUE -- from a `row/*` worktree everything reports as pending PR
disposition, and the test would have passed against the defect it exists to
catch. Suite 47/47, and 5/5 red without the fix.

A STALE MESSAGE IN A TEST. `6603356a` taught `LoadMergedBf16RawNK` to accept
F8_E4M3 shards and rewrote its rejection to name the supported dtypes;
`test_qwen27_dense_forward.cpp:229` still asserted the old "expected BF16", so
`build-test-cpu` and both sanitizer legs failed on it. The expectation now reads
the message the loader raises, and the FP8 merge path that arrived WITHOUT a test
in this file gets one: a mixed BF16+FP8 merged parameter, expectations hand-
computed from E4M3 bytes and the scale (never re-derived through the same
dequant helper the loader calls), plus the per-channel-scale rejection.
7/7, 333 assertions.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-opus-5 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 9, 2026
…message (#210)

Main CI has been red on EVERY merged PR. Two independent defects, both landed
by #178's push (run 31332846716); `cuda-fat-build` in that run was cancelled by
concurrency, not failing.

THE GATE REDDENED MAIN FOR OBEYING IT. `check-role-discipline.py` judges every
commit in the `before..after` range on its own message. A PR landed with a REAL
merge commit pushes the merge AND the branch commits under it: the merge names
the PR, the branch commits were never required to, so each one read as a direct
push. `6603356a` (#178), `e73cbbae` (#204) and `1a02ab4f` (#196) all failed this
way, in both `documentation-checkpoint` and `agent-record`. Arrival is now judged
ONCE, on the commit that lands the change: `merged_pr_content` exempts what a
row/* PR merge brings in. Squash-merges are untouched -- their one commit carries
"(#N)" and passes on its own message. NOT a weakening, and gated as such: only
the SIDE parents count, so `--not parents[0]` keeps a commit pushed straight to
main from being laundered by merging a PR on top, and a merge naming no row and
no PR exempts nothing. Four unit checks build real git history for those cases,
plus the exact `3bbee96e..0cf3dbb` range CI ran, pinned with `has_reached_main`
forced TRUE -- from a `row/*` worktree everything reports as pending PR
disposition, and the test would have passed against the defect it exists to
catch. Suite 47/47, and 5/5 red without the fix.

A STALE MESSAGE IN A TEST. `6603356a` taught `LoadMergedBf16RawNK` to accept
F8_E4M3 shards and rewrote its rejection to name the supported dtypes;
`test_qwen27_dense_forward.cpp:229` still asserted the old "expected BF16", so
`build-test-cpu` and both sanitizer legs failed on it. The expectation now reads
the message the loader raises, and the FP8 merge path that arrived WITHOUT a test
in this file gets one: a mixed BF16+FP8 merged parameter, expectations hand-
computed from E4M3 bytes and the scale (never re-derived through the same
dequant helper the loader calls), plus the per-channel-scale rejection.
7/7, 333 assertions.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-opus-5 [ClaudeCode]

Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants