diff --git a/.agents/NOW.md b/.agents/NOW.md index 4c3239905..bdb9a98d6 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -15,7 +15,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | `SPEC-DSPARK` | **WORKS on 35B**: ON==OFF 48/48. ★fixed engine-wide draft-drop | Draft step ~6x a target step | | State record (#166) | **157 imports = 3,231,342 bytes** at `776c56f1`; 95/95 | Force-update #166; rerun readiness | | Laguna NVFP4 / DeepSeek-V4 decode | **CLOSED, byte-exact, default-ON**: 1.03x vLLM, 1.144x ds4 | Laguna vLLM K-run | -| 27B NVFP4 @`0893e160` | **0.72x -> 0.85x**: FP8 tower native, tokens MATCH, RSS -3.2 GiB | NVFP4 MLP marlin, 68% of roof | +| 27B NVFP4 @`0893e160` | **0.85x**: FP8 native, tokens MATCH; #213 head packed | #213 grid + RSS remeasure post-#150 | | f32-out GEMV audit | **CLAIM WRONG**: 35B runs 41 `CastF32`/step (3.1%) | Fold into the 35B lever | | Invocation-parity | CI guard + checklist landing | build-verify `kGemvHeuristicAlgos` | | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index a3c613a75..416e1d014 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -44,6 +44,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#230](https://github.com/mudler/vllm.cpp/issues/230) | — | `test_agent_record.py`: 7 issue-intake tests never run, and error when they do | bug | | [#222](https://github.com/mudler/vllm.cpp/issues/222) | `SERVE-GATE-ONLINE` | The online-serving gate aborts on current main: harness looks for `examples/server`, the build produces `examples/vllm-server` — paths repaired by `2b262622`/`8fce04d3`; residue is the drift guard, the `27n` key and a model gate that recorded a skip as passed | bug | | [#224](https://github.com/mudler/vllm.cpp/issues/224) | `ENG-DOCS-SITE` | Publish `docs/` as a GitHub Pages site without owning a second copy | feature | +| [#213](https://github.com/mudler/vllm.cpp/issues/213) | `PERF-27B-LMHEAD-FP4` | Qwen3.6 NVFP4 baselines (27B and 35B-A3B) must reach vLLM speed parity | perf | | [#203](https://github.com/mudler/vllm.cpp/issues/203) | `BACKEND-VULKAN` | Vulkan on unified memory holds TWO copies of the weights: 27B peaks at 100.8 GiB RSS and OOM-reboots a Spark | bug | | [#201](https://github.com/mudler/vllm.cpp/issues/201) | `BACKEND-ROCM` | `hipblasGemmEx` overload mismatch in `rocm_matmul_hipblaslt.hip` | bug | | [#199](https://github.com/mudler/vllm.cpp/issues/199) | `BACKEND-METAL-MLX` | macOS MLX build fails on `-Werror` in MLX headers | bug | diff --git a/.agents/specs/perf-27b-lmhead-nvfp4.md b/.agents/specs/perf-27b-lmhead-nvfp4.md new file mode 100644 index 000000000..8bd023368 --- /dev/null +++ b/.agents/specs/perf-27b-lmhead-nvfp4.md @@ -0,0 +1,305 @@ +# PERF-27B-LMHEAD-FP4 — keep the ModelOpt NVFP4 `lm_head` packed + +Issue: [#213](https://github.com/mudler/vllm.cpp/issues/213) +Row: `PERF-27B-LMHEAD-FP4` +Gate model: `nvidia/Qwen3.6-27B-NVFP4` @`0893e1606ff3d5f97a441f405d5fc541a6bdf404` +Base: `origin/main` @`04069bd7`; the round-1 findings were rebased onto +`origin/main` @`723d96a8`, the round-2 findings onto `origin/main` @`a0fa12c7` +(after `ENG-LOAD-DIRECT-UPLOAD` #150). + +## Scope + +The dense Qwen3.6 loader dequantizes a ModelOpt NVFP4 `lm_head` into a BF16 +`[in,out]` operand at load. Keep it packed and run the logits GEMM on the same +Marlin W4A16 family vLLM pins for this head. + +**In scope:** the dense `lm_head` weight path only — loader, weight struct, +resident staging, and the two `DenseLmHead` consumers (gather and non-gather), +the eager `ForwardLogits` arm, and the dense MTP sibling. + +**Out of scope:** the FP8 tower output dtype (`PERF-27B-FP8-BF16-OUT`), the +gate_up merge (`PERF-27B-DENSE-GATEUP-MERGE`), the embedding table, the MoE +path's head, and anything on the `unsloth` repos. + +## The gap, verified against current code + +`LoadLmHeadAnyDtype`'s `U8` branch +(`src/vllm/model_executor/models/qwen3_5_dense_weights.cpp:266-306`) runs +`DequantCtNvfp4WeightToF32` into a full f32 array, rounds to BF16, then +`TransposeBf16` into an `OwnedTensor {in_dim, out_dim}`. The comment at +`:205-211` records this as deliberate deferral: + +> Keeping the head quantized end-to-end would save ~2.3 GiB but needs an +> `lm_head_fp4`-style field on the dense weights plus a forward branch; that is +> a follow-up, not this fix. + +Consequence at decode: the logits GEMM re-reads ~2.543 GB of BF16 every step +where the packed head is ~0.715 GB (`K*N/2` packed + `K*N/16` F8_E4M3 block +scales). Storing the operand transposed to `[K,N]` additionally forces a +row-major NN GEMM, for which no `nvjet_sm121` kernel exists, so cuBLASLt falls +back to the legacy `cutlass_80_tensorop_s16816gemm_bf16_128x64_32x6_nn_align2` — +an Ampere tile on an `sm_121a` part. + +## Upstream anchors + +- `vllm/model_executor/layers/quantization/modelopt.py:2508-2536` — + `ModelOptMixedPrecisionConfig.get_quant_method` accepts `ParallelLMHead` and + returns `ModelOptNvFp4LinearMethod` for `quant_algo == "NVFP4"`. +- `modelopt.py:2491-2496` — `_quantized_layer_prefix_candidates` appends the + bare `lm_head` key, so the mixed scheme is *designed* to resolve a quantized + head. +- `modelopt.py:1249,1283-1284` — `ModelOptNvFp4W4A16LinearMethod` pins + `MarlinNvFp4LinearKernel`, explicitly because the generic priority list would + otherwise first-pick a W4A4 cutlass kernel on this hardware. +- `modelopt.py:1365` — vLLM **deletes** `input_scale` on the W4A16 path + (`process_weights_after_loading`); the placeholder is registered at + `modelopt.py:1358`. Verified against the pinned oracle `555967922`. +- `vllm/model_executor/layers/logits_processor.py:98-133` — `_apply_head` calls + `lm_head.quant_method.apply` every step; nothing materializes BF16. +- `marlin_utils_fp4.py:157-218,221-306` — `prepare_fp4_layer_for_marlin` / + `apply_fp4_marlin_linear`, the repack and apply we already vendored. + +## Design + +1. Add `Nvfp4Weight lm_head_fp4` to the dense weight struct + (`include/vllm/model_executor/models/qwen3_5_dense.h`). +2. In the loader's `U8` branch, stop dequantizing: route through the existing + `IsNvfp4Projection` / `LoadNvfp4AnyNaming` path into `lm_head_fp4`, keeping + the ModelOpt `weight_scale_2`-as-scale convention already handled at + `:274-288`. + + **Deviation, ratified in review:** the design said to retire the stale + "lm_head is never quantized" rule in `IsQwen27QuantizedLinear` + (`qwen3_5_dense_weights.cpp:527` today, not `:506`). It was deliberately NOT + retired: the function has ZERO production callers — only its own routing test + in `test_qwen27_dense_forward.cpp` — so changing it would move no behavior + while invalidating a checked-in expectation. The routing that matters lives in + `LoadDenseLmHead`. +3. A `DenseLogitsF32D` helper gains the packed branch so the tied and `nk` + cases keep one code path. All three consumers route through it: the gathered + and non-gathered paged arms and the eager `ForwardDense` arm. +4. Build the head's resident **pre-capture**, from the registry `prepare` hook + (`Qwen3_5DenseModel::PrepareLmHeadResident`), mirroring the MoE + `PrepareMarlinResident`. `BuildMarlinDenseResident` copies a function-local + host float into the device buffer; built lazily inside a captured region, the + copy source does not outlive the capture. On a backend with NO fp4 GEMM the + same hook builds the dequantized bf16 operand instead, so the fallback never + dequantizes per forward call. +5. `PrepareBf16Resident` needs no head-specific branch: a packed head's bf16 + owner is empty by construction, and the function is only reached under + `IsPlainBf16Qwen3_5Dense`, which is false whenever the head is packed. **The + RSS win comes from the LOADER no longer building the f32 + bf16 arrays**, not + from anything skipped at staging time. +6. The dense MTP ctor (`:6684-6686`) has no `lm_head_fp4_` sibling; add it. + +Gate the whole thing behind `VT_LMHEAD_FP4` (default ON once green) so the A/B +is same-binary and there is an in-binary rollback. + +**BF16 and FP8 heads are untouched.** Those branches keep their current bytes, +so every recorded `unsloth` benchmark is unaffected. + +## Risks + +- **Graph hang, not fault.** Marlin's fp32 reduce spins on lock words; a + workspace that is not zero at allocation hangs forever. Zero at alloc. +- **`IsTrueW4A4()` flip.** This checkpoint ships `lm_head.input_scale`. + Consuming it would select the W4A4 GEMM that vLLM explicitly refuses + (`modelopt.py:1365`) AND make `PrepareLmHeadResident` early-return, silently + skipping the pre-capture build. The rule holds under BOTH spellings: + `LoadCtNvfp4Raw` consumes `input_global_scale` unconditionally (correct for a + 27B TOWER projection, wrong for an output head), so `LoadDenseLmHead` drops the + activation globals unless `VT_MODELOPT_W4A4=1`. Assert + `lm_head_fp4.IsTrueW4A4() == false` for both namings. +- **A compressed-tensors head read as tied.** The model loader's head probe must + accept `.weight_packed` as well as `.weight` + (`DenseCheckpointHasLmHead`); a bare `.weight` probe reads a CT head as + `tie_word_embeddings` and computes the logits off the embedding table. +- **Gate blindness.** `test_qwen27_paged_engine` (235/235) runs `unsloth` + @`890bdef7`, which ships a **BF16** head — that gate cannot see this path. A + fresh greedy continuation on `nvidia`@`0893e160` against the pinned oracle is + mandatory, not optional. + +## Tests + +RED first, `tests/parity/test_qwen27_dense_lmhead_fp4.cpp`: + +1. Synthetic `modelopt_mixed` fixture: `U8 lm_head.weight` + `F8_E4M3 + lm_head.weight_scale` + f32 `lm_head.weight_scale_2`. Assert + `lm_head_fp4.Empty() == false`, the BF16 owner is absent, and the resident + byte count is `K*N/2 + K*N/16`. Red today: the field does not exist. +2. Numerical: logits from the packed head match a reference dequant-then-GEMM + within the Marlin W4A16 tolerance already used by the 32B-NVFP4A16 op tests. +3. Assert `IsTrueW4A4() == false` for the loaded head. + +Added in the review-findings round, each pinned by a mutation that turns it RED: + +4. The same `IsTrueW4A4() == false` assertion under compressed-tensors names + (`weight_packed` / `weight_global_scale` / `input_global_scale`), and + `DenseCheckpointHasLmHead` accepting the CT spelling. +5. Both PAGED lm_head call sites — the gathered (prefill/mixed) and the + non-gathered arm — against the eager reference. The numerical case above runs + only `ForwardDense`, so reverting either paged arm to the bf16 owner was + invisible. +6. The fallback dequant is built ONCE (pointer identity across two forwards), + not per call, and the registry `prepare` hook builds it before any forward. + +Added in the ROUND-2 findings round: + +7. The same case additionally populates an NVFP4 **tower** (every layer's + `mlp.{gate,up,down}_proj_fp4`) and asserts that after two forwards on the same + no-fp4-GEMM backend NONE of them holds a `d_dequant_b`. RED under the mutation + that removes the `keep_dequant_b` guard (6 failing assertions, 3 projections x + 2 layers), which is exactly the round-1 behavior. + +Added in the ROUND-3 findings round: + +8. Case 7 builds its tower with `MakeNvfp4Weight` — direct struct construction — + so it pins what the FORWARD does with a weight that did not opt in and says + nothing about which LOADER may set the flag. No test in the repo called + `LoadDenseMlp`, `LoadDenseAttn` or `LoadQwen3_5DenseWeights` at all. Case 8 + loads a fully-NVFP4 dense layer of BOTH layer types, under BOTH namings + (ModelOpt `weight_scale_2` and compressed-tensors `weight_packed`), through + `LoadQwen3_5DenseLayer` on a synthetic bag, and sweeps every `Nvfp4Weight` the + layer STRUCT owns — fields, not call sites — asserting `keep_dequant_b` is + false on all of them and true on the head. A census assertion + (`populated == 11` per naming: GDN `out_proj` + 3 MLP on the linear-attention + layer, 4 attention + 3 MLP on the full-attention layer) keeps a fixture that + stopped routing NVFP4 from reading as a pass. + + Three mutations, each RED: `keep_dequant_b = true` in `LoadNvfp4AnyNaming`'s + ModelOpt arm (11 failures, the ModelOpt leg), the same in `LoadCtNvfp4Raw` + (11 failures, the compressed-tensors leg), and a genuinely FOURTH setter + outside `LoadNvfp4AnyNaming` — one line after `LoadDenseMlp`'s `down_proj` + load (4 failures, both legs). The first is the reviewer's exact mutation, + under which cases 1-7, `test_qwen27_paged_forward` and `test_mtp_speculator` + all stayed green. + + `LoadQwen3_5DenseLayer` gains a `has`-taking overload in the header. The + resolver-only overload answers `has` with a constant `true`, which forces every + routed projection down the compressed-tensors spelling and so cannot reach the + ModelOpt arm at all. The definition already existed; only the declaration is + new. + +Port anchor: `marlin_utils_fp4.py` tolerances and shapes as used by the existing +NVFP4A16 op tests. + +## Gates + +- Focused: the new test, plus `test_qwen27_paged_engine` 235/235 unchanged. +- Full: CUDA `ctest` on `sm_121a`, clean Release build with + `-DVLLM_CPP_CUTLASS_DIR` and `-DVLLM_CPP_TRITON=ON`. +- Correctness: greedy continuation on `nvidia`@`0893e160`, 32 tokens, + `ignore_eos`, captured from the same warm process as the throughput numbers, + compared against the pinned oracle. Token-exact, or a ratified near-tie with + the oracle's own top-2 margin recorded. +- Speed: `VT_LMHEAD_FP4=0|1` same-binary A/B, one `flock $HOME/gpu.lock`, warm + server, 3 reps per leg, order-alternated, medians of per-rep medians, at + c1/c2/c4/c8. Expected effect is far above the 0.5% noise band, so e2e + throughput resolves it; report `nsys --cuda-graph-trace=node` instance counts + for the logits kernel on both legs as the invocation-parity evidence. +- Memory: peak host RSS on both legs. Expected delta **1.70 GiB**: the bf16 + head is `2*K*N` = 2,543,206,400 B = 2.368 GiB and the packed head is + `K*N/2 + K*N/16` = 715,264,000 B = 0.666 GiB, at the real 248320x5120. The + 21.06 -> 19.36 GiB reading was taken BEFORE #150 rewrote `LoadCtNvfp4Raw` to + borrow mmap'd bytes; that changes what host RSS counts, so the figure is + recorded as OWED a re-measurement rather than carried forward as accepted. + + **That -1.70 GiB is CUDA's, and the sign is not the same everywhere.** CUDA + never dequantizes, so it holds the packed head alone. A backend with no fp4 + GEMM holds the packed head *plus* the one bf16 operand it multiplies against, + 0.666 + 2.368 = 3.034 GiB. Arithmetic from the same K*N, not measured: + + | Backend | before this row | after | delta | + |---|---|---|---| + | CUDA | 2.368 bf16 | 0.666 packed | **-1.70 GiB** | + | Vulkan (unified, #203) | 2.368 host bf16 + 2.368 device copy = 4.736 | 0.666 host packed + 2.368 device bf16 = 3.034 | **-1.70 GiB** | + | plain CPU | 2.368 bf16 | 0.666 + 2.368 = 3.034 | **+0.67 GiB** | + + So #203's backend genuinely improves and plain CPU genuinely regresses by the + packed head's own bytes. The trade is deliberate: CPU pays 0.67 GiB once + instead of rebuilding 2.368 GiB on every decode step, which is what it did + before this row and what round 1 tried to fix for the whole tower at once. + +## Evidence + +`dgx:~/work/vllm.cpp-online-gate/evidence//lmhead-fp4/` — raw A/B legs, +both `nsys` reports, the continuation transcripts from both engines, RSS +samples, and the build recipe. + +## Residency, and the shared-seam exception + +The dequantized bf16 `[K,N]` operand a backend with no fp4 GEMM multiplies +against is kept for the model's lifetime **only for a weight that opts in** +(`Nvfp4Weight::keep_dequant_b`, set by `LoadDenseLmHead` and by nothing else). + +Round 2 found the round-1 fix had cached it for EVERY NVFP4 weight, because the +caching sat inside `MatmulNvfp4F32D` / `MatmulNvfp4Bf16D`, which also serve the +dense MLP, `o_proj`, GDN `out_proj` and the MoE shared experts. `kMatmulNvfp4` is +registered CUDA-only (`cuda_matmul_nvfp4.cu`), so CPU, Vulkan, Metal, HIP and +Tenstorrent all take that fallback: their steady state went from packed-only to +packed plus a bf16 expansion of the WHOLE tower, roughly 4x the packed bytes, on +the backends where issue #203 already reports the 27B peaking at 100.8 GiB. + +Residency is therefore a property of the WEIGHT, not of the GEMM. It is worth its +bytes exactly where one operand is re-read whole every step and there is one of +it — the output head. Alternatives rejected: making it a per-BACKEND switch (the +tower is the problem on every fallback backend, not on a particular one) and +dequantizing the head into `Qwen3_5DenseWeights::lm_head` at prepare time (the +prepare hook may hold BORROWED, const weights, so only `mutable` residency state +is writable there). + +**Shared-seam exception, recorded.** `qwen3_5.cpp` carries a private device +dispatcher (`MatmulNvfp4F32D` / `MatmulNvfp4Bf16D`) parallel to the shared +`dense_nvfp4::MatmulNvfp4W4A16D`. That predates this row: `dense_nvfp4_gemm.h` +was EXTRACTED from `qwen3_5.cpp`'s anonymous namespace and its own SCOPE comment +records that the true-W4A4 (fp4-activation) path stays private to `qwen3_5.cpp`, +so the two also carry independent `Dev`, `MakeTensor`, `ResidentNvfp4` and +`DequantNvfp4ToBLayout` copies. Unifying them is a refactor this row does not do. +What this row does instead is put the opt-in on the SHARED data type +(`Nvfp4Weight`, `qwen3_5_weights.h`), which both dispatchers can read, and +default it OFF. + +**They CAN still disagree, and the claim is only that they do not.** The shared +fallback at `dense_nvfp4_gemm.h:626-631` ignores `keep_dequant_b` outright: it +rebuilds the `K*N` bf16 operand per call for every weight, opted in or not. So a +weight that opted in and then reached `MatmulNvfp4W4A16D` would get the per-call +temporary rather than the resident. That does not happen today because no weight +reachable from `MatmulNvfp4W4A16D` opts in, and if it ever did the divergence is +in the benign direction: the shared seam UNDER-caches, never over-caches, so it +cannot reintroduce the whole-tower expansion this finding was about. Honoring the +flag in the shared seam belongs to the unification refactor, not here. The PR +body and commit message claim only the dense `lm_head`, never "every fp4 +projection". + +**The opt-in has exactly ONE setter, and a test says so.** `LoadDenseLmHead` is +it. `LoadNvfp4AnyNaming` is the single function every dense NVFP4 TOWER +projection flows through (MLP gate/up/down via `LoadDenseMlp`, attention q/k/v/o +via `LoadAttnDense`, GDN `out_proj` via `LoadGdnDense`), so one added line there +reopens the round-2 blocker verbatim — and the CUDA gate cannot see it, because +`kMatmulNvfp4` is registered on CUDA and `ResidentNvfp4DequantB` is never +reached. Round 3 found that every gate stayed green under exactly that mutation. +Test 8 below closes it. + +**A backend-asymmetric hard abort, deliberate.** `ResidentNvfp4DequantB` opens +with `VT_CHECK(w.keep_dequant_b, ...)`, so `PrepareLmHeadResident` +(`qwen3_5.cpp:6534`) ABORTS for any `Qwen3_5DenseWeights` whose `lm_head_fp4` is +non-empty but did not come from `LoadDenseLmHead` — and only on a backend with no +fp4 GEMM, because CUDA returns at `:6533` first. It is unreachable in production +(the loader is the only producer of a packed head), and it is preferred to a +silent fall-through because a head that quietly lost its opt-in would rebuild +2.54 GB per step with no symptom but throughput. A test that hand-builds an +`lm_head_fp4` is the one caller that can trip it: +`tests/vllm/models/test_qwen27_paged_forward.cpp:1377` does exactly that and +passes only because it never calls `Prepare`. + +## Stop conditions + +- Stop and report `NEEDS_DECISION` if the packed head cannot be made + token-exact-or-ratified against the oracle. +- Stop if the head's `weight_scale_2` convention does not match the ModelOpt + reading already used at `:274-288`. +- Do not widen scope into the FP8 tower, the gate_up merge, or the MoE head. + +## Outcome + +Pending. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 2162fea61..de451c689 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -94,6 +94,20 @@ the same metric at higher concurrency (c8 p99 ITL 0.86x, but 1.055x at c16 and | Roof | 20.42 GiB over about 273 GB/s: vLLM's 79 ms/token is about 95% of the bandwidth limit, ours 96 ms/token about 78% | | | | | Peak host RSS | 21.0 GiB, down from 24.2 GiB, because the FP8 tower is no longer expanded to BF16 | | | | +#### NVFP4 `lm_head` kept packed (`PERF-27B-LMHEAD-FP4`, #213) + +| Axis | Packed (`VT_LMHEAD_FP4=1`) | Dequant (`=0`) | Result | +|---|---:|---:|---| +| Peak host RSS | 19.36 GiB | 21.06 GiB | **-1.70 GiB**, but measured BEFORE `ENG-LOAD-DIRECT-UPLOAD` (#150) made `LoadCtNvfp4Raw` borrow mmap'd bytes, which moves the RSS accounting; re-measurement OWED | +| Peak host RSS, non-CUDA | | | **Arithmetic, not measured.** A backend with no fp4 GEMM keeps packed + one bf16 operand: Vulkan **-1.70 GiB** (it used to stage a host bf16 head *and* a device copy), plain CPU **+0.67 GiB**. See `docs/USAGE.md` | +| Greedy continuation | identical to the dequant leg, byte for byte | | SOLID | +| `test_qwen27_paged_engine` | 235/235 | 235/235 | unchanged | +| tok/s, leg A / leg B | 11.197 / 11.193 | 9.418 / 10.163 | **INDICATIVE ONLY** | +| Reading, throughput | packed faster in all four legs; packed legs agree to 0.04%, dequant legs disagree by 7.9% | | DIRECTION established, MAGNITUDE not | +| Owed | binding grid: 3 reps per leg, order-alternated, c1/c2/c4/c8, medians of per-rep medians, before any ratio is quoted | | PENDING | +| Method | same model and revision as the table above, same-binary A/B on GB10 | | | +| Provenance | the row's fresh reviewer, gate checkpoint `nvidia/Qwen3.6-27B-NVFP4`@`0893e1606ff3d5f97a441f405d5fc541a6bdf404` | | | + ### Qwen3.6-35B-A3B by concurrency | Concurrency | 1 | 2 | 4 | 8 | 16 | 32 | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 0c58b01bf..21ad01a86 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -90,6 +90,7 @@ portable/reference path. In normal operation leave them unset. | `VT_CONV_REG` | on (CUDA GDN) | The non-register-tiled short causal convolution | | `VT_CONV_EXACT_CHUNKS` | on (CUDA GDN prefill) | Use `=0` for the legacy sequence-serial causal-conv mapping; default mirrors vLLM's exact `(sequence, 8-token chunk)` descriptor and is byte-identical | | `VT_MODELOPT_W4A4` | `0` (Qwen3.6 dense ModelOpt NVFP4) | ModelOpt NVFP4 checkpoints ship a per-tensor `input_scale` next to every projection. Consuming it sets `Nvfp4Weight::alpha`, which flips `IsTrueW4A4()` and routes the weight to the fp4-ACTIVATION GEMM; on `nvidia/Qwen3.6-27B-NVFP4` that produced incoherent text, so the default leaves `alpha` at 0 and takes the W4A16 weight-only dispatcher (verified coherent). Set `1` to consume `input_scale` and take the W4A4 path | +| `VT_LMHEAD_FP4` | **on** (Qwen3.6 dense NVFP4 `lm_head`) | Keeps a ModelOpt/compressed-tensors NVFP4 output head PACKED (`Qwen3_5DenseWeights::lm_head_fp4`) so the logits GEMM reads `K*N/2 + K*N/16` bytes per step instead of the `2*K*N` of a dequantized bf16 operand (~0.715 GB vs ~2.543 GB on `nvidia/Qwen3.6-27B-NVFP4`), and the operand keeps its on-disk `[N,K]` orientation instead of forcing the row-major NN GEMM that has no `nvjet_sm121` kernel. Mirrors vLLM, which resolves a quantized `lm_head` through `ModelOptNvFp4W4A16LinearMethod` (`modelopt.py:2491-2496,2508-2536`) and never materializes bf16 (`logits_processor.py:98-133`). Measured peak host RSS 21.06 -> 19.36 GiB (**-1.70 GiB**: the bf16 head is 2,543,206,400 B = 2.368 GiB, the packed head 715,264,000 B = 0.666 GiB) — measured before #150 changed the RSS accounting, so the figure is owed a re-measurement. On a backend with NO fp4 GEMM the head additionally keeps ONE dequantized `2*K*N` bf16 operand, built at prepare time; no other NVFP4 projection keeps one, so a quantized tower is never expanded (issue #203). The head is W4A16 under BOTH spellings: the on-disk activation divisor (`input_scale` / `input_global_scale`) is dropped for the head unless `VT_MODELOPT_W4A4=1`, because vLLM's `ModelOptNvFp4W4A16LinearMethod` deletes it (`modelopt.py:1365`). `=0` is the same-binary rollback to dequantize-at-load. BF16, FP8, GGUF and tied heads are unaffected either way (row `PERF-27B-LMHEAD-FP4`, issue #213) | | `VT_FA2_PREFILL` | on (CUDA) | The portable prefill attention instead of the vendored FA2 | | `VT_FA2_DECODE` | on (CUDA) | The portable decode attention instead of the vendored FA2 | | `VT_FA2_DECODE_4B` | on (CUDA, Qwen3.5-4B) | The portable paged decode attention instead of the ratio-4 vendored FA2 path; the 27B and 35B selectors are unchanged | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index a4e61967f..f42e23883 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -65,6 +65,7 @@ are our reading of their documented behavior, not measurements. | Format | vllm.cpp | vLLM | SGLang | llama.cpp | |---|---|---|---|---| | NVFP4 (W4A4 and W4A16 Marlin) | ✅ | ✅ | ✅ | ☐ | +| NVFP4 `lm_head` kept packed (no dequant at load) | ✅ `VT_LMHEAD_FP4` default-ON, #213; CUDA-gated on `nvidia/Qwen3.6-27B-NVFP4`@`0893e160` (continuations byte-identical packed vs dequant, `test_qwen27_paged_engine` 235/235; RSS -1.70 GiB owed a re-measure) | ✅ | ☐ | ☐ | | GGUF k-quants and i-quants | ✅ (CPU grouped keep-quant MoE took a bf16-activation regression in `b4f5610a`; found by bisect and fixed 2026-08-06) | ☐ | ☐ | ✅ | | AWQ | ◐ CPU dequant | ✅ | ✅ | ☐ | | GPTQ | ◐ CPU dequant | ✅ | ✅ | ☐ | diff --git a/docs/STATUS.md b/docs/STATUS.md index 71a06b08e..bfe352897 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -75,7 +75,7 @@ token-for-token correctness against the pinned oracle. | Capability | State | Notes | |---|---|---| -| Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | +| Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` **0.85x BEHIND**; its NVFP4 `lm_head` PACKED (#213; RSS remeasure due) | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | | Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact GDN chunks default ON and byte-identical to rollback. Local A/B: total/output +2.152%, TTFT -2.945%, TPOT/ITL -1.920%; sealed-vLLM comparison 1.021x throughput, 1.086x TTFT, 1.025x TPOT, +233 MiB VRAM ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | diff --git a/docs/USAGE.md b/docs/USAGE.md index 0d78d159a..f4b4f0af3 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -285,6 +285,34 @@ independently selectable with `VT_CPU_Q8_DOT`, `VT_CPU_QUANT_MMLA`, and while an unavailable forced tier fails closed. The exact accepted values are listed in [ENVIRONMENT.md](ENVIRONMENT.md). +### The NVFP4 output head + +On a Qwen3.6 dense checkpoint whose `lm_head` is stored NVFP4 (ModelOpt +`weight`/`weight_scale`/`weight_scale_2`, or compressed-tensors +`weight_packed`/`weight_global_scale`) the head is kept **packed** and the logits +GEMM runs on it directly, as vLLM does. Nothing is dequantized at load, so the +head costs `K*N/2 + K*N/16` bytes instead of `2*K*N`, about 0.715 GB instead of +2.543 GB on `nvidia/Qwen3.6-27B-NVFP4` (measured peak host RSS 21.06 to 19.36 +GiB, a 1.70 GiB saving on CUDA; the figure is owed a re-measurement after +`ENG-LOAD-DIRECT-UPLOAD` changed the RSS accounting). + +That accounting is CUDA's. A backend with no fp4 GEMM (CPU, Vulkan, Metal, HIP, +Tenstorrent) has to multiply against a dequantized bf16 copy, so on those the +head costs the packed bytes **plus** one `2*K*N` operand, built once when the +model is prepared rather than per call — 0.666 + 2.368 = 3.034 GiB on the same +checkpoint. The sign of the change therefore depends on the backend: on Vulkan, +which used to stage a host bf16 head *and* a device copy of it, the head goes +4.736 to 3.034 GiB, the same **-1.70 GiB**; on plain CPU it goes 2.368 to 3.034, +a **+0.67 GiB** regression, paid once instead of rebuilding 2.368 GiB on every +decode step as that backend did before. Only the head is kept that way; every +other NVFP4 projection dequantizes per call, so a quantized tower is never +expanded in memory. The head runs W4A16 under both namings: the on-disk +activation divisor next to it (`input_scale`, or `input_global_scale` in the +compressed-tensors spelling) is NOT consumed unless `VT_MODELOPT_W4A4=1`, +matching vLLM, which deletes it on this path. Set `VT_LMHEAD_FP4=0` for a +same-binary A/B that restores the old dequantize-at-load owner. BF16, FP8, GGUF +and `tie_word_embeddings` heads are unaffected by either setting. + ### Validating a staged release archive Release verification reads only a freshly extracted archive, never files from diff --git a/include/vllm/model_executor/models/qwen3_5_dense.h b/include/vllm/model_executor/models/qwen3_5_dense.h index 1aada1c09..6b6279623 100644 --- a/include/vllm/model_executor/models/qwen3_5_dense.h +++ b/include/vllm/model_executor/models/qwen3_5_dense.h @@ -97,14 +97,24 @@ struct Qwen3_5DenseLayerWeights { DenseMlpWeights mlp; // every layer has a dense MLP }; -// Whole dense-model text weights. `lm_head` is always materialized bf16 here, but -// the CHECKPOINT may store it BF16, FP8 (per-channel scale) or ModelOpt NVFP4 — -// the 27B NVFP4 publishers disagree, and revisions of one repo disagree with each -// other (issue #164). LoadLmHeadAnyDtype dequantizes all three to this operand. +// Whole dense-model text weights. The CHECKPOINT may store the head BF16, FP8 +// (per-channel scale) or ModelOpt NVFP4 — the 27B NVFP4 publishers disagree, and +// revisions of one repo disagree with each other (issue #164). BF16/FP8 are +// materialized into `lm_head`, NVFP4 stays PACKED in `lm_head_fp4` +// (PERF-27B-LMHEAD-FP4, issue #213); exactly one is populated. struct Qwen3_5DenseWeights { OwnedTensor embed_tokens; // bf16 [vocab, H] (NOT transposed; embed lookup) OwnedTensor final_norm; // bf16 [H] OwnedTensor lm_head; // bf16 [H, vocab] (dequantized -> Matmul-B layout) + // NVFP4-resident output head [N=vocab, K=H], kept in the on-disk orientation the + // fp4 GEMMs read. Mirrors Qwen3_5MoeWeights::lm_head_fp4 and vLLM's own decision + // to leave the head quantized: get_quant_method accepts ParallelLMHead + // (modelopt.py:2508-2536) over the bare `lm_head` key (modelopt.py:2491-2496), + // so ModelOptNvFp4W4A16LinearMethod — pinning MarlinNvFp4LinearKernel + // (modelopt.py:1249,1283-1284) — resolves it and logits_processor._apply_head + // calls quant_method.apply every step (logits_processor.py:98-133). Empty on + // every BF16/FP8/GGUF/tied checkpoint. + Nvfp4Weight lm_head_fp4; // Mirrors tie_word_embeddings: logits reuse embed_tokens as raw [V,H] // torch-Linear storage, so no second host/device owner is created. bool tied_lm_head = false; @@ -135,6 +145,25 @@ OwnedTensor LoadLmHeadAnyDtype(const TensorResolver& get, const std::function& has, const std::string& name); +// PERF-27B-LMHEAD-FP4 (issue #213). Load the dense output head into EXACTLY ONE +// of the two owners: a ModelOpt/compressed-tensors NVFP4 head stays PACKED in +// `fp4_out`, every other storage form is materialized bf16 [in, out] into +// `bf16_out` by LoadLmHeadAnyDtype. `proj` omits the trailing ".weight". +void LoadDenseLmHead(const TensorResolver& get, + const std::function& has, + const std::string& proj, OwnedTensor& bf16_out, + Nvfp4Weight& fp4_out); + +// True when the checkpoint ships an EXPLICIT head under either naming +// (`.weight`, or `.weight_packed` for compressed-tensors NVFP4); +// false means `tie_word_embeddings`. +bool DenseCheckpointHasLmHead(const std::function& has, + const std::string& proj); + +// VT_LMHEAD_FP4 (default ON): the in-binary rollback for the packed head. `0` +// restores the dequantize-at-load owner, so the A/B is same-binary. +bool DenseLmHeadFp4Enabled(); + Fp8Weight LoadFp8RawShared(const TensorResolver& get, const std::string& proj); OwnedTensor MaterializeCtNvfp4Bf16Transposed(const TensorResolver& get, @@ -161,6 +190,17 @@ Qwen3_5DenseLayerWeights LoadQwen3_5DenseLayer(const TensorResolver& get, const std::string& layer_type, int64_t layer_idx); +// The same load with an EXPLICIT presence probe — what `LoadQwen3_5Dense` calls +// per layer. The resolver-only overload above answers `has` with a constant +// `true`, which forces every routed projection down the compressed-tensors +// spelling; a checkpoint that mixes namings (the ModelOpt `weight_scale_2` form, +// or an FP8/BF16 projection next to an NVFP4 one) needs the real probe. Exposed +// so the loader gate can drive a whole synthetic layer through the SAME routing +// production takes. +Qwen3_5DenseLayerWeights LoadQwen3_5DenseLayer( + const TensorResolver& get, const std::function& has, + const std::string& layer_type, int64_t layer_idx); + // Full dense-model load across the given shards. Uses config.num_hidden_layers // and config.layer_types. Text path only — the vision tower (model.visual.*) // and image/video merger are DEFERRED (notes §0.1). The checkpoint's MTP @@ -188,6 +228,13 @@ class Qwen3_5DenseModel { static void PrepareBf16Resident(const Qwen3_5DenseWeights& weights, vt::Queue& queue); + // PERF-27B-LMHEAD-FP4 (issue #213). Build the resident form of the packed + // `lm_head_fp4` THIS backend's logits GEMM consumes: the Marlin W4A16 repack + // on CUDA (PRE-CAPTURE), else the dequantized bf16 [K,N] operand. Inert when + // the head is not packed. Called from the registry `prepare` hook. + static void PrepareLmHeadResident(const Qwen3_5DenseWeights& weights, + vt::Queue& queue); + // Batched PAGED dense forward — the 27B analogue of Qwen3_5Model::Forward. // Same signature/structure (paged KV cache for the full-attn layers, batched // GDN recurrent state for the GDN layers, the f32 residual thread), reusing the diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h index bc5bac507..022ce8d67 100644 --- a/include/vllm/model_executor/models/qwen3_5_weights.h +++ b/include/vllm/model_executor/models/qwen3_5_weights.h @@ -238,6 +238,16 @@ struct Nvfp4Weight { // path. Uploaded once from the persistent `alpha` member; the diagnostic host // scalar path leaves this null. mutable std::shared_ptr d_alpha; + // OPT-IN lifetime residency for the DEQUANTIZED bf16 [K=in, N=out] Matmul-B + // operand the backends with NO fp4 GEMM multiply against (CPU / Vulkan / Metal / + // HIP / Tenstorrent; CUDA never dequantizes). Default OFF, and it must stay a + // per-WEIGHT opt-in: the operand is a bf16 expansion of ~4x the packed bytes, so + // holding one per tower projection is the double-residency that OOM-reboots a + // Spark on Vulkan (#203). The dense loader opts in the OUTPUT HEAD alone — one + // weight, re-read whole every step (~2.54 GB a step rebuilt per call at the + // 27B's 248320x5120); everything else keeps a per-call copy. + bool keep_dequant_b = false; + mutable std::shared_ptr d_dequant_b; // Resident Marlin constants (issue #237; see ResidentSlot). `resident_marlin` // is this weight's own repack; `resident_marlin_pair` is the fused gate+up diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 2d491187f..287f4cdb9 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -1187,6 +1187,39 @@ std::vector DequantNvfp4ToBLayout(const Nvfp4Weight& w) { return io; } +// The SAME bf16 [K=in, N=out] operand, uploaded ONCE and kept resident on the +// weight (mirror of ResidentNvfp4, same Backend deleter). OPT-IN per weight +// (`keep_dequant_b`, qwen3_5_weights.h): only the output head is worth a lifetime +// bf16 expansion of ~4x its packed bytes. +Tensor ResidentNvfp4DequantB(Dev d, const Nvfp4Weight& w) { + VT_CHECK(w.keep_dequant_b, "nvfp4: dequant-B residency is opt-in per weight"); + if (!w.d_dequant_b) { + const std::vector wb = DequantNvfp4ToBLayout(w); + const size_t nb = wb.size() * sizeof(uint16_t); + void* p = d.b.Alloc(nb); + d.b.Copy(d.q, p, wb.data(), nb); + Backend* bk = &d.b; + w.d_dequant_b = std::shared_ptr(p, [bk](void* q) { bk->Free(q); }); + } + return MakeTensor(w.d_dequant_b.get(), DType::kBF16, d.q.device, {w.k, w.n}); +} + +// out[M,N] = x[M,K] @ dequant(w).T — the fallback both device dispatchers take on +// a backend with NO fp4 GEMM (CPU registers only kMatmulNvfp4Fp4; Vulkan/Metal +// neither kMatmulNvfp4 nor the Marlin grouped GEMM). A weight that did not opt in +// keeps the PER-CALL temporary it has always had; caching the whole NVFP4 tower +// would quadruple its steady-state bytes on exactly those backends. +void MatmulNvfp4DequantB(Dev d, Tensor& out, const Tensor& x, + const Nvfp4Weight& w) { + if (w.keep_dequant_b) { + vt::Matmul(d.q, out, x, ResidentNvfp4DequantB(d, w)); + return; + } + const std::vector wb = DequantNvfp4ToBLayout(w); + DBuf dwb(d, DType::kBF16, {w.k, w.n}, wb.data()); + vt::Matmul(d.q, out, x, dwb.t()); +} + // y[M,N] f32 = x[M,K] bf16 @ dequant(w).T, w fp4-resident [N=out, K=in]. Drops // in for MatmulF32 where the weight is NVFP4 (experts/shared/lm_head). std::vector MatmulNvfp4F32(Dev d, const std::vector& x, int64_t M, @@ -2549,7 +2582,7 @@ DBuf SharedGateUpFusedMarlinD(Dev d, const Tensor& x, const Nvfp4Weight& gw, #endif // VT_MARLIN_NVFP4 DBuf MatmulNvfp4F32D(Dev d, const Tensor& x, const Nvfp4Weight& w) { - const int64_t M = x.shape[0], K = x.shape[1], N = w.n; + const int64_t M = x.shape[0], N = w.n; if (vllm::platforms::GetPlatform(d.q.device.type).cutlass_fp4_supported() && w.IsTrueW4A4() && TrueW4A4Enabled()) return MatmulNvfp4Fp4D(d, x, w, DType::kF32); #ifdef VT_MARLIN_NVFP4 @@ -2565,18 +2598,33 @@ DBuf MatmulNvfp4F32D(Dev d, const Tensor& x, const Nvfp4Weight& w) { Nvfp4Dev dw = ResidentNvfp4(d, w); vt::MatmulNvfp4(d.q, dout.t(), x, dw.packed, dw.scale, w.scale2); } else { - std::vector wb = DequantNvfp4ToBLayout(w); - DBuf dwb(d, DType::kBF16, {K, N}, wb.data()); - vt::Matmul(d.q, dout.t(), x, dwb.t()); + MatmulNvfp4DequantB(d, dout.t(), x, w); } return dout; } +// The ONE dense-gate logits GEMM: y[M,vocab] f32 = x[M,H] @ lm_head. +// PERF-27B-LMHEAD-FP4 (issue #213). A ModelOpt NVFP4 head stays PACKED, so the +// GEMM reads K*N/2 + K*N/16 bytes per step instead of the 2*K*N of a dequantized +// bf16 operand (~0.715 GB vs ~2.543 GB at the real 248320x5120), and keeps its +// on-disk [N,K] orientation instead of forcing the row-major NN GEMM that has no +// nvjet_sm121 kernel. Mirrors logits_processor._apply_head -> +// lm_head.quant_method.apply (logits_processor.py:98-133). Every dense consumer +// (eager ForwardDense, the gathered and non-gathered paged arms) routes here, so +// exactly one head layout is selected; the bf16 arm keeps both of its shapes. +DBuf DenseLogitsF32D(Dev d, const Tensor& x, const Qwen3_5DenseWeights& weights) { + if (!weights.lm_head_fp4.Empty()) + return MatmulNvfp4F32D(d, x, weights.lm_head_fp4); + const OwnedTensor& lm_head = DenseLmHead(weights); + return lm_head.nk ? MatmulBf16LogitsF32D(d, x, lm_head) + : MatmulF32D(d, x, lm_head); +} + // Same as MatmulNvfp4F32D but bf16 output (the down/o/out_proj sinks that feed // the residual add). CUDA: fp4-resident vt::MatmulNvfp4 (bf16 out). CPU: the // DequantNvfp4ToBLayout fallback (no CPU MatmulNvfp4 kernel). DBuf MatmulNvfp4Bf16D(Dev d, const Tensor& x, const Nvfp4Weight& w) { - const int64_t M = x.shape[0], K = x.shape[1], N = w.n; + const int64_t M = x.shape[0], N = w.n; if (vllm::platforms::GetPlatform(d.q.device.type).cutlass_fp4_supported() && w.IsTrueW4A4() && TrueW4A4Enabled()) return MatmulNvfp4Fp4D(d, x, w, DType::kBF16); #ifdef VT_MARLIN_NVFP4 @@ -2589,9 +2637,7 @@ DBuf MatmulNvfp4Bf16D(Dev d, const Tensor& x, const Nvfp4Weight& w) { Nvfp4Dev dw = ResidentNvfp4(d, w); vt::MatmulNvfp4(d.q, dout.t(), x, dw.packed, dw.scale, w.scale2); } else { - std::vector wb = DequantNvfp4ToBLayout(w); - DBuf dwb(d, DType::kBF16, {K, N}, wb.data()); - vt::Matmul(d.q, dout.t(), x, dwb.t()); + MatmulNvfp4DequantB(d, dout.t(), x, w); } return dout; } @@ -6456,6 +6502,43 @@ void Qwen3_5Model::PrepareMarlinResident(const Qwen3_5MoeWeights& weights, #endif } +// PERF-27B-LMHEAD-FP4 (issue #213). Build whatever resident form of the PACKED +// dense head THIS backend's logits GEMM will actually consume, once, at prepare +// time. Inert on every BF16/FP8/GGUF/tied head (`lm_head_fp4` empty). +// +// CUDA/Marlin: prepare time is strictly BEFORE any decode-graph capture, and that +// matters — BuildMarlinDenseResident Allocs, launches the repack, and Copies a +// host float whose source is a function-local temporary. Same arm as +// Qwen3_5Model::PrepareMarlinResident's lm_head build above. A backend with NO +// fp4 GEMM builds the dequantized bf16 [K,N] operand here instead, so the head — +// the one weight that opted into it — never pays it on the forward path. +void Qwen3_5DenseModel::PrepareLmHeadResident(const Qwen3_5DenseWeights& weights, + vt::Queue& queue) { + if (weights.lm_head_fp4.Empty()) return; + Dev d{vt::GetBackend(queue.device.type), queue}; +#ifdef VT_MARLIN_NVFP4 + // Build under EXACTLY the guard MatmulNvfp4F32D uses to select the Marlin + // GEMM, so a configuration that will not take that path never builds for it. + if (!weights.lm_head_fp4.IsTrueW4A4() && MarlinMoeEnabled() && + vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, queue.device.type)) { + BuildMarlinDenseResident(d, weights.lm_head_fp4, + MarlinDenseResidentFor(&weights.lm_head_fp4)); + d.b.Synchronize(d.q); + return; + } +#endif + // Same selection order as MatmulNvfp4F32D: the fp4-activation and packed-GEMM + // arms stage the packed bytes lazily and NOT per call, so only the + // dequantizing fallback needs eager work here. + if (vllm::platforms::GetPlatform(queue.device.type).cutlass_fp4_supported() && + weights.lm_head_fp4.IsTrueW4A4() && TrueW4A4Enabled()) { + return; + } + if (vt::OpRegistered(vt::OpId::kMatmulNvfp4, queue.device.type)) return; + (void)ResidentNvfp4DequantB(d, weights.lm_head_fp4); + d.b.Synchronize(d.q); +} + void Qwen3_5DenseModel::PrepareBf16Resident( const Qwen3_5DenseWeights& weights, vt::Queue& queue) { VT_CHECK(platforms::GetPlatform(queue.device.type).needs_weight_staging(), @@ -6474,6 +6557,9 @@ void Qwen3_5DenseModel::PrepareBf16Resident( raw(weights.embed_tokens); raw(weights.final_norm); + // PERF-27B-LMHEAD-FP4: already a no-op for a PACKED head (empty bf16 owner, and + // IsPlainBf16Qwen3_5Dense is false whenever the head is packed); its resident is + // built by PrepareLmHeadResident. raw(DenseLmHead(weights)); for (const Qwen3_5DenseLayerWeights& layer : weights.layers) { raw(layer.input_layernorm); @@ -6689,10 +6775,9 @@ std::vector Qwen3_5DenseModel::ForwardDense( DBuf dnorm(d, DType::kBF16, {T, H}); vt::RmsNorm(d.q, dnorm.t(), hidden.t(), dfn, vt::RmsNormArgs{eps, true}, &res.t()); - // lm_head is unquantized bf16 in the 27B (notes §3.6): the one host Download. - const OwnedTensor& lm_head = DenseLmHead(weights); - DBuf dlogits = lm_head.nk ? MatmulBf16LogitsF32D(d, dnorm.t(), lm_head) - : MatmulF32D(d, dnorm.t(), lm_head); + // lm_head (the one host Download): PACKED NVFP4 (PERF-27B-LMHEAD-FP4) when the + // checkpoint ships a ModelOpt/CT NVFP4 head, else the bf16/tied owner. + DBuf dlogits = DenseLogitsF32D(d, dnorm.t(), weights); std::vector logits(static_cast(T) * vocab); dlogits.Download(d, logits.data()); return logits; @@ -6704,7 +6789,10 @@ Qwen3_5MTPModel::Qwen3_5MTPModel(const Qwen3_5MTPWeights& weights, : weights_(&weights), config_(&config), embed_tokens_(&target.embed_tokens), - lm_head_(&DenseLmHead(target)) { + lm_head_(&DenseLmHead(target)), + // PERF-27B-LMHEAD-FP4: the drafter shares the TARGET's head, so it must see + // the packed one too. Empty on every BF16/FP8/tied dense target. + lm_head_fp4_(&target.lm_head_fp4) { VT_CHECK(weights.kind == Qwen3_5MTPKind::kDense, "qwen3_5 MTP: dense target requires dense MTP weights"); } @@ -7034,21 +7122,17 @@ static DBuf DenseForwardLayers(Dev d, const Tensor& hidden_in, } // Logits gather-before-lm_head (prefill/mixed): same semantics as the 35B path. - // lm_head is unquantized bf16 in the 27B (notes §3.6). Pure-decode / graph - // replay pass empty indices (identity) → the full [T,vocab] path. + // Both arms route through DenseLogitsF32D (PERF-27B-LMHEAD-FP4). Pure-decode / + // graph replay pass empty indices (identity) → the full [T,vocab] path. const bool do_gather = !logits_indices.empty() && static_cast(logits_indices.size()) < T; if (do_gather) { const int64_t n_out = static_cast(logits_indices.size()); DBuf dgather(d, DType::kBF16, {n_out, H}); GatherRows(d, dgather.ptr(), dnorm.t(), logits_indices, H); - const OwnedTensor& lm_head = DenseLmHead(weights); - return lm_head.nk ? MatmulBf16LogitsF32D(d, dgather.t(), lm_head) - : MatmulF32D(d, dgather.t(), lm_head); + return DenseLogitsF32D(d, dgather.t(), weights); } - const OwnedTensor& lm_head = DenseLmHead(weights); - return lm_head.nk ? MatmulBf16LogitsF32D(d, dnorm.t(), lm_head) - : MatmulF32D(d, dnorm.t(), lm_head); + return DenseLogitsF32D(d, dnorm.t(), weights); } // Full eager dense paged forward body: embed (host token_ids) then the capturable diff --git a/src/vllm/model_executor/models/qwen3_5_dense.cpp b/src/vllm/model_executor/models/qwen3_5_dense.cpp index 2b67bb142..d212738d8 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense.cpp @@ -102,9 +102,12 @@ std::unique_ptr LoadQwen3_5DenseModel( void PrepareQwen3_5Dense(LoadedModel& model, const HfConfig& config, vt::Queue& queue) { - (void)model; (void)config; - (void)queue; + // PERF-27B-LMHEAD-FP4 (issue #213): build the packed lm_head's resident HERE — + // on CUDA before the runner captures a decode graph, elsewhere before the first + // forward pays the dequant. Inert on every BF16/FP8/GGUF/tied checkpoint. + auto& qwen = static_cast(model); + Qwen3_5DenseModel::PrepareLmHeadResident(qwen.weights(), queue); } ForwardLogits ForwardQwen3_5Dense(LoadedModel& model, diff --git a/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp index ed12efae8..dc904bf3d 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp @@ -142,6 +142,14 @@ void StageAndReleaseLoadedDense(Qwen3_5DenseWeights& weights, (void)ReleaseResidentQwen3_5DenseHostWeights(weights); } +// VT_MODELOPT_W4A4 (default 0): consume a projection's on-disk activation +// divisor, setting `alpha` and so flipping `IsTrueW4A4()` to the fp4-ACTIVATION +// GEMM (docs/ENVIRONMENT.md). ONE reader, so the spellings cannot drift. +bool ModelOptW4A4OptIn() { + const char* w4a4 = std::getenv("VT_MODELOPT_W4A4"); + return w4a4 != nullptr && w4a4[0] == '1'; +} + // One compressed-tensors NVFP4 W4A4 Linear -> RAW fp4-resident Nvfp4Weight kept // in the on-disk [N=out, K=in] orientation vt::MatmulNvfp4 reads directly (notes // §5 step-6a — the throughput path; NO bf16 materialization). Reads the CT @@ -219,11 +227,11 @@ Nvfp4Weight LoadCtNvfp4Raw(const TensorResolver& get, const std::string& proj) { // head to FP8 with a PER-OUTPUT-CHANNEL scale, and nvidia/Qwen3.6-27B-NVFP4 ships // a ModelOpt NVFP4 head; both hit the old unconditional BF16 assert. // -// All three land on the SAME bf16 [in, out] Matmul-B operand the logits GEMM -// already consumes, so the forward is untouched and a BF16 head stays byte-exact -// (identical call, no dequant). Keeping the head quantized end-to-end would save -// ~2.3 GiB but needs an `lm_head_fp4`-style field on the dense weights plus a -// forward branch; that is a follow-up, not this fix. +// BF16 and FP8 heads land on the SAME bf16 [in, out] Matmul-B operand the logits +// GEMM already consumes, so a BF16 head stays byte-exact. The NVFP4 form no +// longer reaches this function: LoadDenseLmHead routes it to the PACKED +// `lm_head_fp4` (PERF-27B-LMHEAD-FP4, issue #213), and the U8 branch below +// survives ONLY as the VT_LMHEAD_FP4=0 in-binary rollback. // // ModelOpt vs compressed-tensors global-scale convention: CT stores the value as // a DIVISOR and `DequantCtNvfp4WeightToF32` reciprocates it internally, whereas @@ -377,8 +385,8 @@ Nvfp4Weight LoadNvfp4AnyNaming(const TensorResolver& get, const TensorExists& ha // A/B(VT_MODELOPT_W4A4=1): default W4A16. ModelOpt ships `input_scale` on every // projection, but consuming it flips IsTrueW4A4() and routes to the // fp4-activation GEMM; leaving alpha at 0 keeps the weight-only dispatcher. - const char* w4a4 = std::getenv("VT_MODELOPT_W4A4"); - if (w4a4 != nullptr && w4a4[0] == '1' && has(proj + ".input_scale")) { + const bool w4a4_opt_in = ModelOptW4A4OptIn(); + if (w4a4_opt_in && has(proj + ".input_scale")) { const float is = ReadF32Scalar(get(proj + ".input_scale")); if (is != 0.0F) { r.input_global_scale_inv = is; @@ -499,6 +507,51 @@ DenseMlpWeights LoadDenseMlp(const TensorResolver& get, const TensorExists& has, } // namespace +bool DenseLmHeadFp4Enabled() { + const char* v = std::getenv("VT_LMHEAD_FP4"); + return v == nullptr || v[0] != '0'; +} + +void LoadDenseLmHead(const TensorResolver& get, const TensorExists& has, + const std::string& proj, OwnedTensor& bf16_out, + Nvfp4Weight& fp4_out) { + fp4_out = Nvfp4Weight{}; + bf16_out = OwnedTensor{}; + // PERF-27B-LMHEAD-FP4 (issue #213). An NVFP4 head stays PACKED, through the SAME + // LoadNvfp4AnyNaming every other NVFP4 projection takes, so the ModelOpt vs + // compressed-tensors global-scale convention is handled in exactly one place. + // vLLM's mixed scheme likewise resolves a quantized head (modelopt.py:2491-2496). + if (DenseLmHeadFp4Enabled() && IsNvfp4Projection(has, proj)) { + fp4_out = LoadNvfp4AnyNaming(get, has, proj); + // The head is W4A16, whatever the naming. `LoadNvfp4AnyNaming` decides + // activation-quant per SPELLING — the ModelOpt arm ignores `input_scale` + // unless VT_MODELOPT_W4A4=1, but `LoadCtNvfp4Raw` consumes + // `input_global_scale` UNCONDITIONALLY, correct for a TOWER projection of the + // 27B compressed-tensors checkpoint, which really is W4A4. An output head is + // not one: vLLM resolves it through ModelOptNvFp4W4A16LinearMethod, which + // DELETES input_scale (modelopt.py:1365; registered at :1358) and pins + // MarlinNvFp4LinearKernel (modelopt.py:1249,1283-1284). A set alpha would (a) + // take the fp4-activation GEMM vLLM refuses here and (b) make + // PrepareLmHeadResident early-return, skipping the pre-capture Marlin build. + if (!ModelOptW4A4OptIn()) { + fp4_out.input_global_scale_inv = 0.0F; + fp4_out.alpha = 0.0F; + } + // The ONE weight that opts into a lifetime dequant-B resident where there is + // no fp4 GEMM (qwen3_5_weights.h): the head is re-read whole every step. + fp4_out.keep_dequant_b = true; + return; + } + bf16_out = LoadLmHeadAnyDtype(get, has, proj + ".weight"); +} + +bool DenseCheckpointHasLmHead(const TensorExists& has, const std::string& proj) { + // A bare `.weight` probe misses a compressed-tensors head, whose only + // weight tensor is `.weight_packed`; no head at all means + // `tie_word_embeddings`, so missing CT ties the logits to the embedding table. + return has(proj + ".weight") || IsNvfp4Projection(has, proj); +} + OwnedTensor LoadMergedBf16RawNK(const TensorResolver& get, const std::vector& names) { // Extracted to the shared dense_weight_loaders.h (SEAM GAP #3); this retains @@ -632,8 +685,8 @@ Qwen3_5DenseWeights LoadQwen3_5Dense(const std::vector& shards, LoadModelBf16Direct(get, "model.language_model.norm.weight"); // The 27B owns an explicit head; smaller Qwen3.5 checkpoints tie logits to // the embedding table and omit lm_head.weight. - if (has("lm_head.weight")) { - w.lm_head = LoadLmHeadAnyDtype(get, has, "lm_head.weight"); + if (DenseCheckpointHasLmHead(has, "lm_head")) { + LoadDenseLmHead(get, has, "lm_head", w.lm_head, w.lm_head_fp4); } else { w.tied_lm_head = true; w.embed_tokens.nk = true; @@ -652,6 +705,9 @@ Qwen3_5DenseWeights LoadQwen3_5Dense(const std::vector& shards, } bool IsPlainBf16Qwen3_5Dense(const Qwen3_5DenseWeights& weights) { + // A PACKED head (PERF-27B-LMHEAD-FP4) is not plain bf16: this staging path + // only knows how to stage OwnedTensors. + if (!weights.lm_head_fp4.Empty()) return false; for (const Qwen3_5DenseLayerWeights& layer : weights.layers) { if (!layer.mlp.gate_proj_fp4.Empty() || !layer.mlp.up_proj_fp4.Empty() || !layer.mlp.down_proj_fp4.Empty()) { diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 296e66b3f..a47b9d9fb 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1244,6 +1244,12 @@ target_compile_definitions(test_qwen27_paged_engine PRIVATE target_include_directories(test_qwen27_paged_engine PRIVATE ${CMAKE_SOURCE_DIR}/tests/parity ${CMAKE_SOURCE_DIR}/src) +# `PERF-27B-LMHEAD-FP4` (issue #213): the ModelOpt NVFP4 lm_head stays PACKED. +# Synthetic (no checkpoint, no GPU) BECAUSE test_qwen27_paged_engine runs the +# unsloth @890bdef7 snapshot, whose head is BF16 — that gate cannot see this path. +vllm_cpp_add_test(test_qwen27_dense_lmhead_fp4 + parity/test_qwen27_dense_lmhead_fp4.cpp) + # SPEC-DFLASH D2 — the DFlash draft-forward parity gate vs the dumped vLLM DFlash # draft reference (fc combine + context-free block forward per-stage hidden + # proposed ids). GPU + z-lab-draft + NVFP4-target + fixture gated (skips on CPU/CI). diff --git a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp new file mode 100644 index 000000000..8808f83f7 --- /dev/null +++ b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp @@ -0,0 +1,706 @@ +// PERF-27B-LMHEAD-FP4 (issue #213) — keep the ModelOpt NVFP4 `lm_head` PACKED. +// `nvidia/Qwen3.6-27B-NVFP4` ships a ModelOpt NVFP4 output head (`lm_head.weight` +// U8 + `lm_head.weight_scale` F8_E4M3 + `lm_head.weight_scale_2` f32). The dense +// loader used to DEQUANTIZE it into a bf16 [in,out] Matmul-B owner, so the logits +// GEMM re-read ~2.543 GB every decode step where the packed head is ~0.715 GB. +// vLLM keeps it quantized instead (anchors on `Qwen3_5DenseWeights::lm_head_fp4`, +// qwen3_5_dense.h). +// +// These cases pin the LOADER ROUTING, the NUMERICS and the RESIDENCY of that +// decision. They are synthetic (no checkpoint, no GPU) on purpose: the 235/235 +// `test_qwen27_paged_engine` gate runs `unsloth`@890bdef7, whose head is BF16, so +// that gate is BLIND to this path. +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vllm/model_executor/layers/quantization/compressed_tensors/nvfp4_emulation.h" +#include "vllm/model_executor/model_loader/nvfp4_dequant.h" +#include "vllm/model_executor/model_loader/safetensors_reader.h" +#include "vllm/model_executor/models/model_registry.h" +#include "vllm/model_executor/models/qwen3_5_dense.h" +#include "vllm/transformers_utils/hf_config.h" +#include "vt/backend.h" +#include "vt/dtype.h" + +using vllm::DenseCheckpointHasLmHead; +using vllm::HfConfig; +using vllm::LoadDenseLmHead; +using vllm::ModelRegistry; +using vllm::Nvfp4Weight; +using vllm::OwnedTensor; +using vllm::Qwen3_5DenseLayerWeights; +using vllm::Qwen3_5DenseModel; +using vllm::Qwen3_5DenseWeights; +using vllm::StTensor; +using vt::DType; + +namespace { + +// --- an in-memory safetensors stand-in (mirrors test_qwen3_5_lm_head_dtypes) --- + +struct Fake { + std::string dtype; + std::vector shape; + std::vector bytes; +}; + +class Bag { + public: + void Put(const std::string& name, Fake f) { items_[name] = std::move(f); } + + vllm::TensorResolver Resolver() { + return [this](const std::string& name) -> const StTensor& { + auto it = items_.find(name); + REQUIRE_MESSAGE(it != items_.end(), "missing tensor: " << name); + Fake& f = it->second; + StTensor& v = views_[name]; + v.dtype = f.dtype; + v.shape = f.shape; + v.data = f.bytes.data(); + v.nbytes = f.bytes.size(); + return v; + }; + } + + std::function Has() { + return [this](const std::string& n) { return items_.count(n) != 0; }; + } + + private: + std::unordered_map items_; + std::unordered_map views_; +}; + +uint64_t Mix(uint64_t x) { + x += 0x9E3779B97F4A7C15ULL; + x = (x ^ (x >> 30)) * 0xBF58476D1CE4E5B9ULL; + x = (x ^ (x >> 27)) * 0x94D049BB133111EBULL; + return x ^ (x >> 31); +} + +float RandV(uint64_t seed) { + const double u = static_cast(Mix(seed) >> 40) / static_cast(1 << 24); + return static_cast(u * 0.16 - 0.08); +} + +// One deterministic ModelOpt NVFP4 fixture. The nibbles and group scale bytes are +// chosen DIRECTLY (exact E2M1 codes, exact powers-of-two fp8 block scales) so the +// value a correct dequant must produce is exact and known here: +// w[r][c] = sign * kE2M1Lut[idx] * F8E4M3ToF32(block_scale) * weight_scale_2 +// `weight_scale_2` is the ModelOpt convention: the SCALE, not the CT divisor. +constexpr float kFixtureScale2 = 0.125F; + +struct Nvfp4Fixture { + std::vector packed; // [N, K/2] + std::vector scale; // [N, K/16] + std::vector dequant; // [N, K] the exact expected value +}; + +Nvfp4Fixture MakeNvfp4Fixture(int64_t n, int64_t k, uint64_t seed) { + Nvfp4Fixture f; + f.packed.assign(static_cast(n) * static_cast(k / 2), 0); + f.scale.assign(static_cast(n) * static_cast(k / 16), 0); + f.dequant.assign(static_cast(n) * static_cast(k), 0.0F); + // Exact e4m3 powers of two: 0.25, 0.5, 1.0, 2.0. + const float kBlockScales[4] = {0.25F, 0.5F, 1.0F, 2.0F}; + for (int64_t r = 0; r < n; ++r) { + for (int64_t g = 0; g < k / 16; ++g) { + const float bs = + kBlockScales[Mix(seed + static_cast(r * 977 + g)) & 3U]; + f.scale[static_cast(r) * static_cast(k / 16) + + static_cast(g)] = vllm::F32ToF8E4M3(bs); + for (int64_t j = 0; j < 16; ++j) { + const int64_t c = g * 16 + j; + const uint64_t h = Mix(seed + static_cast(r * 131071 + c)); + const int idx = static_cast(h % 8U); + const bool neg = ((h >> 8) & 1U) != 0U; + const float mag = vllm::kE2M1Lut[idx]; + const uint8_t nib = vllm::Fp4ToNibble(neg ? -mag : mag); + const size_t byte = static_cast(r) * static_cast(k / 2) + + static_cast(c / 2); + // low nibble first (the on-disk E2M1 packing both dequants read). + if (c % 2 == 0) + f.packed[byte] = static_cast(f.packed[byte] | (nib & 0x0FU)); + else + f.packed[byte] = static_cast(f.packed[byte] | (nib << 4)); + f.dequant[static_cast(r) * static_cast(k) + + static_cast(c)] = + (neg ? -mag : mag) * bs * kFixtureScale2; + } + } + } + return f; +} + +void PutModelOptHead(Bag& bag, const std::string& proj, int64_t n, int64_t k, + const Nvfp4Fixture& f, bool with_input_scale) { + bag.Put(proj + ".weight", Fake{"U8", {n, k / 2}, f.packed}); + bag.Put(proj + ".weight_scale", Fake{"F8_E4M3", {n, k / 16}, f.scale}); + Fake s2{"F32", {}, std::vector(4)}; + const float v = kFixtureScale2; + std::memcpy(s2.bytes.data(), &v, 4); + bag.Put(proj + ".weight_scale_2", std::move(s2)); + if (with_input_scale) { + // The gate checkpoint DOES ship `lm_head.input_scale`. Consuming it would flip + // IsTrueW4A4() and select the W4A4 GEMM vLLM refuses here (modelopt.py:1365). + Fake is{"F32", {}, std::vector(4)}; + const float iv = 0.0625F; + std::memcpy(is.bytes.data(), &iv, 4); + bag.Put(proj + ".input_scale", std::move(is)); + } +} + +// The SAME fp4 bytes under compressed-tensors names: `weight_global_scale` is a +// DIVISOR (1/scale), and the per-Linear `input_global_scale` shipped next to it +// is exactly the activation divisor that must NOT be consumed on an output head. +void PutCtNvfp4Head(Bag& bag, const std::string& proj, int64_t n, int64_t k, + const Nvfp4Fixture& f) { + bag.Put(proj + ".weight_packed", Fake{"U8", {n, k / 2}, f.packed}); + bag.Put(proj + ".weight_scale", Fake{"F8_E4M3", {n, k / 16}, f.scale}); + Fake wgs{"F32", {}, std::vector(4)}; + const float divisor = 1.0F / kFixtureScale2; // CT stores the RECIPROCAL + std::memcpy(wgs.bytes.data(), &divisor, 4); + bag.Put(proj + ".weight_global_scale", std::move(wgs)); + Fake igs{"F32", {}, std::vector(4)}; + const float iv = 16.0F; + std::memcpy(igs.bytes.data(), &iv, 4); + bag.Put(proj + ".input_global_scale", std::move(igs)); +} + +// The fixture bytes as an NVFP4 TOWER projection [N=out, K=in]. +Nvfp4Weight MakeNvfp4Weight(int64_t n, int64_t k, uint64_t seed) { + const Nvfp4Fixture f = MakeNvfp4Fixture(n, k, seed); + Nvfp4Weight w; + w.n = n; + w.k = k; + w.scale2 = kFixtureScale2; + w.packed.dtype = w.scale.dtype = DType::kI8; + w.packed.rank = w.scale.rank = 2; + w.packed.shape[0] = w.scale.shape[0] = n; + w.packed.shape[1] = k / 2; + w.scale.shape[1] = k / 16; + w.packed.bytes.assign(f.packed.begin(), f.packed.end()); + w.scale.bytes.assign(f.scale.begin(), f.scale.end()); + return w; +} + +// One TOWER projection [N=out, K=in] under either spelling. `LoadNvfp4AnyNaming` +// has an arm for each — `LoadCtNvfp4Raw` for compressed-tensors, its own body for +// ModelOpt — so a residency opt-in added to one is invisible to a fixture that +// only ever exercises the other. +void PutNvfp4Proj(Bag& bag, const std::string& proj, int64_t n, int64_t k, + uint64_t seed, bool ct_naming) { + const Nvfp4Fixture f = MakeNvfp4Fixture(n, k, seed); + if (ct_naming) { + PutCtNvfp4Head(bag, proj, n, k, f); + } else { + // ModelOpt ships an `input_scale` next to EVERY projection; a tower + // projection is where the loader may legally ignore it at the default + // `VT_MODELOPT_W4A4=0`. + PutModelOptHead(bag, proj, n, k, f, /*with_input_scale=*/true); + } +} + +uint16_t F32ToBf16(float v) { + uint32_t bits = 0; + std::memcpy(&bits, &v, sizeof(bits)); + const uint32_t lsb = (bits >> 16) & 1U; + bits += 0x7FFFU + lsb; + return static_cast(bits >> 16); +} + +float Bf16ToF32(uint16_t h) { + const uint32_t bits = static_cast(h) << 16; + float v = 0.0F; + std::memcpy(&v, &bits, sizeof(v)); + return v; +} + +Fake MakeBf16(const std::vector& shape, uint64_t seed) { + int64_t total = 1; + for (int64_t s : shape) total *= s; + Fake f{"BF16", shape, std::vector(static_cast(total) * 2)}; + for (int64_t i = 0; i < total; ++i) { + const uint16_t h = F32ToBf16(RandV(seed + static_cast(i))); + std::memcpy(f.bytes.data() + static_cast(i) * 2, &h, 2); + } + return f; +} + +// One whole synthetic dense decoder layer with EVERY routed projection stored +// NVFP4 (ModelOpt spelling) and everything else BF16, under the exact tensor +// names `LoadQwen3_5DenseLayer` resolves. `linear_attention` quantizes the GDN +// `out_proj`, `full_attention` quantizes q/k/v/o_proj, and both quantize the MLP +// — which is every `LoadNvfp4AnyNaming` call site in the dense loader. +void PutNvfp4DenseLayer(Bag& bag, const HfConfig& c, const std::string& type, + int64_t idx, bool ct_naming) { + const std::string base = + "model.language_model.layers." + std::to_string(idx) + "."; + const uint64_t s = 3000 + static_cast(idx) * 700; + const int64_t H = c.hidden_size, I = c.intermediate_size; + const int64_t Hq = c.num_attention_heads, Hkv = c.num_key_value_heads, + Dh = c.head_dim; + const int64_t Hk = c.linear_num_key_heads, Hv = c.linear_num_value_heads, + Dk = c.linear_key_head_dim, Dv = c.linear_value_head_dim, + Kw = c.linear_conv_kernel_dim; + const int64_t key_dim = Hk * Dk, value_dim = Hv * Dv, + conv_dim = 2 * key_dim + value_dim; + bag.Put(base + "input_layernorm.weight", MakeBf16({H}, s + 1)); + bag.Put(base + "post_attention_layernorm.weight", MakeBf16({H}, s + 2)); + if (type == "linear_attention") { + const std::string la = base + "linear_attn."; + // The in-projections are on vLLM's `ignore` list and stay BF16 in the on-disk + // torch-Linear [out, in] orientation; only `out_proj` is quantized. + bag.Put(la + "in_proj_qkv.weight", MakeBf16({conv_dim, H}, s + 10)); + bag.Put(la + "in_proj_z.weight", MakeBf16({value_dim, H}, s + 20)); + bag.Put(la + "in_proj_b.weight", MakeBf16({Hv, H}, s + 30)); + bag.Put(la + "in_proj_a.weight", MakeBf16({Hv, H}, s + 40)); + bag.Put(la + "conv1d.weight", MakeBf16({conv_dim, 1, Kw}, s + 50)); + bag.Put(la + "A_log", MakeBf16({Hv}, s + 60)); + bag.Put(la + "dt_bias", MakeBf16({Hv}, s + 70)); + bag.Put(la + "norm.weight", MakeBf16({Dv}, s + 80)); + PutNvfp4Proj(bag, la + "out_proj", H, value_dim, s + 90, ct_naming); + } else { + const std::string sa = base + "self_attn."; + PutNvfp4Proj(bag, sa + "q_proj", Hq * Dh, H, s + 110, ct_naming); + PutNvfp4Proj(bag, sa + "k_proj", Hkv * Dh, H, s + 120, ct_naming); + PutNvfp4Proj(bag, sa + "v_proj", Hkv * Dh, H, s + 130, ct_naming); + PutNvfp4Proj(bag, sa + "o_proj", H, Hq * Dh, s + 140, ct_naming); + bag.Put(sa + "q_norm.weight", MakeBf16({Dh}, s + 150)); + bag.Put(sa + "k_norm.weight", MakeBf16({Dh}, s + 160)); + } + const std::string mlp = base + "mlp."; + PutNvfp4Proj(bag, mlp + "gate_proj", I, H, s + 210, ct_naming); + PutNvfp4Proj(bag, mlp + "up_proj", I, H, s + 220, ct_naming); + PutNvfp4Proj(bag, mlp + "down_proj", H, I, s + 230, ct_naming); +} + +// EVERY `Nvfp4Weight` a loaded dense layer can own, by struct field rather than +// by loader call site — which is what makes the sweep below survive a loader +// path that does not exist yet. +std::vector> LayerNvfp4Weights( + const Qwen3_5DenseLayerWeights& lw) { + return { + {"linear_attn.out_proj", &lw.gdn.out_proj_fp4}, + {"self_attn.q_proj", &lw.attn.q_proj_fp4}, + {"self_attn.k_proj", &lw.attn.k_proj_fp4}, + {"self_attn.v_proj", &lw.attn.v_proj_fp4}, + {"self_attn.o_proj", &lw.attn.o_proj_fp4}, + {"mlp.gate_proj", &lw.mlp.gate_proj_fp4}, + {"mlp.up_proj", &lw.mlp.up_proj_fp4}, + {"mlp.down_proj", &lw.mlp.down_proj_fp4}, + }; +} + +// --- the small synthetic dense model (shape scaffold from +// test_qwen27_dense_forward.cpp; hidden/vocab widened to Marlin-shaped) --- + +OwnedTensor MakeOwned(DType dt, std::vector shape, uint64_t seed) { + OwnedTensor t; + t.dtype = dt; + t.rank = static_cast(shape.size()); + int64_t n = 1; + for (int i = 0; i < t.rank; ++i) { + t.shape[i] = shape[static_cast(i)]; + n *= shape[static_cast(i)]; + } + if (dt == DType::kBF16) { + t.bytes.resize(static_cast(n) * 2); + auto* p = reinterpret_cast(t.bytes.data()); + for (int64_t i = 0; i < n; ++i) + p[i] = vt::F32ToBF16(RandV(seed + static_cast(i))); + } else { + t.bytes.resize(static_cast(n) * 4); + auto* p = reinterpret_cast(t.bytes.data()); + for (int64_t i = 0; i < n; ++i) p[i] = RandV(seed + static_cast(i)); + } + return t; +} + +HfConfig MakeConfig() { + HfConfig c; + c.model_type = "qwen3_5_text"; + c.architectures = {"Qwen3_5ForConditionalGeneration"}; + c.hidden_size = 128; // == the head's K + c.num_hidden_layers = 2; + c.vocab_size = 256; // == the head's N + c.num_attention_heads = 6; + c.num_key_value_heads = 2; + c.head_dim = 8; + c.layer_types = {"linear_attention", "full_attention"}; + c.intermediate_size = 16; + c.num_experts = 0; + c.linear_num_key_heads = 2; + c.linear_num_value_heads = 6; + c.linear_key_head_dim = 8; + c.linear_value_head_dim = 8; + c.linear_conv_kernel_dim = 4; + c.rope_theta = 10000.0; + c.rotary_dim = 4; + c.rms_norm_eps = 1e-6; + c.max_position_embeddings = 64; + return c; +} + +Qwen3_5DenseWeights MakeWeights(const HfConfig& c) { + Qwen3_5DenseWeights w; + const int64_t H = c.hidden_size, V = c.vocab_size; + const int64_t Hq = c.num_attention_heads, Hkv = c.num_key_value_heads, + Dh = c.head_dim; + const int64_t Hk = c.linear_num_key_heads, Hv = c.linear_num_value_heads, + Dk = c.linear_key_head_dim, Dv = c.linear_value_head_dim, + Kw = c.linear_conv_kernel_dim; + const int64_t key_dim = Hk * Dk, value_dim = Hv * Dv, + conv_dim = 2 * key_dim + value_dim; + w.embed_tokens = MakeOwned(DType::kBF16, {V, H}, 11); + w.final_norm = MakeOwned(DType::kBF16, {H}, 12); + for (int64_t l = 0; l < c.num_hidden_layers; ++l) { + const uint64_t s = 1000 + static_cast(l) * 5000; + Qwen3_5DenseLayerWeights lw; + lw.is_linear_attention = + (c.layer_types[static_cast(l)] == "linear_attention"); + lw.input_layernorm = MakeOwned(DType::kBF16, {H}, s + 1); + lw.post_attention_layernorm = MakeOwned(DType::kBF16, {H}, s + 2); + if (lw.is_linear_attention) { + lw.gdn.in_proj_qkv = MakeOwned(DType::kBF16, {H, conv_dim}, s + 10); + lw.gdn.in_proj_z = MakeOwned(DType::kBF16, {H, value_dim}, s + 20); + lw.gdn.in_proj_b = MakeOwned(DType::kBF16, {H, Hv}, s + 30); + lw.gdn.in_proj_a = MakeOwned(DType::kBF16, {H, Hv}, s + 40); + lw.gdn.conv1d_weight = MakeOwned(DType::kBF16, {conv_dim, Kw}, s + 50); + lw.gdn.a_log = MakeOwned(DType::kF32, {Hv}, s + 60); + lw.gdn.dt_bias = MakeOwned(DType::kF32, {Hv}, s + 70); + lw.gdn.norm_weight = MakeOwned(DType::kBF16, {Dv}, s + 80); + lw.gdn.out_proj = MakeOwned(DType::kBF16, {value_dim, H}, s + 90); + } else { + lw.attn.q_proj = MakeOwned(DType::kBF16, {H, 2 * Hq * Dh}, s + 10); + lw.attn.k_proj = MakeOwned(DType::kBF16, {H, Hkv * Dh}, s + 20); + lw.attn.v_proj = MakeOwned(DType::kBF16, {H, Hkv * Dh}, s + 30); + lw.attn.o_proj = MakeOwned(DType::kBF16, {Hq * Dh, H}, s + 40); + lw.attn.q_norm = MakeOwned(DType::kBF16, {Dh}, s + 50); + lw.attn.k_norm = MakeOwned(DType::kBF16, {Dh}, s + 60); + } + const int64_t I = c.intermediate_size; + lw.mlp.gate_proj = MakeOwned(DType::kBF16, {H, I}, s + 501); + lw.mlp.up_proj = MakeOwned(DType::kBF16, {H, I}, s + 502); + lw.mlp.down_proj = MakeOwned(DType::kBF16, {I, H}, s + 503); + w.layers.push_back(std::move(lw)); + } + return w; +} + +vt::Queue CpuQ() { + return vt::Queue{vt::Device{vt::DeviceType::kCPU, 0}, nullptr}; +} + +// The bf16 [K,N] Matmul-B owner the OLD loader produced for this head: dequant, +// round to bf16, transpose. Built from the fixture's own exact values, NOT by +// calling either production dequant, so the numerical case below compares against +// an INDEPENDENT reference. +OwnedTensor ReferenceBf16Head(const Nvfp4Fixture& f, int64_t n, int64_t k) { + OwnedTensor o; + o.dtype = DType::kBF16; + o.rank = 2; + o.shape[0] = k; + o.shape[1] = n; + o.bytes.resize(static_cast(n) * static_cast(k) * 2); + auto* p = reinterpret_cast(o.bytes.data()); + for (int64_t r = 0; r < n; ++r) + for (int64_t c = 0; c < k; ++c) + p[static_cast(c) * static_cast(n) + static_cast(r)] = + F32ToBf16(f.dequant[static_cast(r) * static_cast(k) + + static_cast(c)]); + return o; +} + +} // namespace + +// ── 1. The loader keeps the ModelOpt NVFP4 head PACKED ─────────────────────── +TEST_CASE("qwen27 dense lm_head: a ModelOpt NVFP4 head stays PACKED (no bf16 owner)") { + constexpr int64_t N = 256, K = 128; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 7); + Bag bag; + PutModelOptHead(bag, "lm_head", N, K, f, /*with_input_scale=*/true); + + OwnedTensor bf16; + Nvfp4Weight fp4; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", bf16, fp4); + + REQUIRE_FALSE(fp4.Empty()); + // The point: NOTHING is materialized to bf16 (the old U8 branch built a [K,N] + // bf16 owner, ~2.543 GB at the real 248320x5120). + CHECK(bf16.Empty()); + CHECK(fp4.n == N); + CHECK(fp4.k == K); + // Resident bytes: K*N/2 packed E2M1 + K*N/16 fp8-e4m3 block scales. + CHECK(fp4.packed.bytes.size() == static_cast(K) * static_cast(N) / 2); + CHECK(fp4.scale.bytes.size() == static_cast(K) * static_cast(N) / 16); + // ModelOpt's weight_scale_2 IS the scale (not the CT divisor). + CHECK(fp4.scale2 == doctest::Approx(kFixtureScale2)); + // ── 3. and the `lm_head.input_scale` this fixture ships must NOT flip the head + // to W4A4: vLLM DELETES it on the W4A16 head path (modelopt.py:1365), and + // consuming it would route the head to the fp4-activation GEMM. + CHECK_FALSE(fp4.IsTrueW4A4()); + CHECK(fp4.alpha == 0.0F); +} + +// ── 3b. The SAME rule under compressed-tensors names ───────────────────────── +// `LoadCtNvfp4Raw` consumes `input_global_scale` UNCONDITIONALLY, correct for a +// TOWER projection of the 27B CT checkpoint, which really is W4A4. An output head +// is not one (modelopt.py:1365): a W4A4 head would take the fp4-activation GEMM +// AND skip the pre-capture Marlin build, which early-returns on IsTrueW4A4(). +TEST_CASE("qwen27 dense lm_head: a compressed-tensors NVFP4 head is W4A16, not W4A4") { + constexpr int64_t N = 256, K = 128; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 23); + Bag bag; + PutCtNvfp4Head(bag, "lm_head", N, K, f); + + // A CT head's ONLY weight tensor is `lm_head.weight_packed`, so a bare + // `lm_head.weight` probe reads it as `tie_word_embeddings` and computes the + // logits off the embedding table. + CHECK(DenseCheckpointHasLmHead(bag.Has(), "lm_head")); + Bag tied; // a genuinely tied checkpoint ships no head under either name + tied.Put("model.language_model.embed_tokens.weight", MakeBf16({8, 4}, 5)); + CHECK_FALSE(DenseCheckpointHasLmHead(tied.Has(), "lm_head")); + + OwnedTensor bf16; + Nvfp4Weight fp4; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", bf16, fp4); + + REQUIRE_FALSE(fp4.Empty()); + CHECK(bf16.Empty()); + CHECK(fp4.n == N); + CHECK(fp4.k == K); + // CT stores the global scale as a DIVISOR; scale2 is its reciprocal. + CHECK(fp4.scale2 == doctest::Approx(kFixtureScale2)); + // The head must land on the SAME W4A16 dispatcher the ModelOpt spelling takes. + CHECK(fp4.alpha == 0.0F); + CHECK_FALSE(fp4.IsTrueW4A4()); +} + +TEST_CASE("qwen27 dense lm_head: a BF16 head is unaffected (the benchmarked form)") { + Bag bag; + bag.Put("lm_head.weight", MakeBf16({8, 4}, 3)); + + OwnedTensor bf16; + Nvfp4Weight fp4; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", bf16, fp4); + + CHECK(fp4.Empty()); + REQUIRE_FALSE(bf16.Empty()); + REQUIRE(bf16.rank == 2); + CHECK(bf16.shape[0] == 4); // in + CHECK(bf16.shape[1] == 8); // out +} + +TEST_CASE("qwen27 dense lm_head: VT_LMHEAD_FP4=0 restores the dequantized owner") { + constexpr int64_t N = 256, K = 128; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 13); + Bag bag; + PutModelOptHead(bag, "lm_head", N, K, f, /*with_input_scale=*/true); + +#ifdef _WIN32 + _putenv_s("VT_LMHEAD_FP4", "0"); +#else + setenv("VT_LMHEAD_FP4", "0", 1); +#endif + OwnedTensor bf16; + Nvfp4Weight fp4; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", bf16, fp4); +#ifdef _WIN32 + _putenv_s("VT_LMHEAD_FP4", ""); +#else + unsetenv("VT_LMHEAD_FP4"); +#endif + + CHECK(fp4.Empty()); + REQUIRE_FALSE(bf16.Empty()); + CHECK(bf16.shape[0] == K); + CHECK(bf16.shape[1] == N); + // The in-binary rollback must reproduce the OLD dequant exactly. + const auto* p = reinterpret_cast(bf16.bytes.data()); + for (int64_t r = 0; r < N; r += 37) + for (int64_t c = 0; c < K; c += 11) + CHECK(Bf16ToF32(p[static_cast(c) * static_cast(N) + + static_cast(r)]) == + doctest::Approx(f.dequant[static_cast(r) * static_cast(K) + + static_cast(c)])); +} + +// ── 2. The packed head computes the SAME logits as the dequantized head ────── +TEST_CASE("qwen27 dense lm_head: packed-head logits match the dequant-then-GEMM reference") { + const HfConfig c = MakeConfig(); + const int64_t N = c.vocab_size, K = c.hidden_size; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 17); + Bag bag; + PutModelOptHead(bag, "lm_head", N, K, f, /*with_input_scale=*/true); + + OwnedTensor unused_bf16; + Nvfp4Weight fp4; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", unused_bf16, fp4); + REQUIRE_FALSE(fp4.Empty()); + + const std::vector ids{3, 11, 40, 200}; + const std::vector pos{0, 1, 2, 3}; + vt::Queue q = CpuQ(); + + // Reference arm: the bf16 [K,N] owner the OLD loader produced. + Qwen3_5DenseWeights ref = MakeWeights(c); + ref.lm_head = ReferenceBf16Head(f, N, K); + const std::vector want = + Qwen3_5DenseModel::ForwardDense(ids, pos, ref, c, q); + + // Packed arm: identical model, head kept packed. + Qwen3_5DenseWeights got_w = MakeWeights(c); + got_w.lm_head_fp4 = fp4; + const std::vector got = + Qwen3_5DenseModel::ForwardDense(ids, pos, got_w, c, q); + + REQUIRE(got.size() == want.size()); + REQUIRE(want.size() == ids.size() * static_cast(N)); + double max_abs = 0.0; + double scale = 1e-6; + for (size_t i = 0; i < want.size(); ++i) { + REQUIRE(std::isfinite(got[i])); + max_abs = std::max(max_abs, std::fabs(static_cast(got[i] - want[i]))); + scale = std::max(scale, std::fabs(static_cast(want[i]))); + } + // Both arms consume the SAME fp4 codes; the only divergence allowed is the + // reference arm's bf16 rounding plus GEMM accumulation order (Marlin W4A16 + // tolerance band). + CHECK(max_abs / scale < 2e-2); +} + +// ── 5. The packed head's resident is built ONCE, at PREPARE time, for the HEAD +// ── ALONE ──────────────────────────────────────────────────────────────────── +// Backends with no fp4 GEMM (CPU here; Vulkan and Metal register neither +// kMatmulNvfp4 nor the Marlin grouped GEMM) fall back to `vt::Matmul` on a +// dequantized bf16 [K,N] operand. Three ways to lose that, none visible to a +// numerical assertion: rebuild it inside the GEMM (2.54 GB per decode step at the +// real 248320x5120); drop the PrepareLmHeadResident call so the forward builds it +// (on CUDA that same call is the PRE-CAPTURE Marlin build); or keep one for EVERY +// NVFP4 weight, quadrupling the tower's steady-state bytes on these same backends +// — the double-residency of issue #203. +TEST_CASE("qwen27 dense lm_head: the packed head's resident is built once, at prepare") { + const HfConfig c = MakeConfig(); + const int64_t N = c.vocab_size, K = c.hidden_size; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 37); + Bag bag; + PutModelOptHead(bag, "lm_head", N, K, f, /*with_input_scale=*/true); + + OwnedTensor unused_bf16; + Nvfp4Weight fp4; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", unused_bf16, fp4); + REQUIRE_FALSE(fp4.Empty()); + + Qwen3_5DenseWeights w = MakeWeights(c); + w.lm_head_fp4 = fp4; + // An NVFP4 TOWER as well as an NVFP4 head, both on the same fallback. + const int64_t I = c.intermediate_size; + for (Qwen3_5DenseLayerWeights& lw : w.layers) { + lw.mlp.gate_proj_fp4 = MakeNvfp4Weight(I, K, 41); + lw.mlp.up_proj_fp4 = MakeNvfp4Weight(I, K, 43); + lw.mlp.down_proj_fp4 = MakeNvfp4Weight(K, I, 47); + } + vt::Queue q = CpuQ(); + + // The registry `prepare` hook builds it BEFORE any forward runs. + const std::unique_ptr model = + vllm::BorrowQwen3_5DenseLoadedModel(w); + REQUIRE(w.lm_head_fp4.d_dequant_b == nullptr); + ModelRegistry::Prepare(*model, c, q); + REQUIRE(w.lm_head_fp4.d_dequant_b != nullptr); + const void* first = w.lm_head_fp4.d_dequant_b.get(); + + // ...and no forward rebuilds it. + const std::vector ids{3, 11}, pos{0, 1}; + (void)Qwen3_5DenseModel::ForwardDense(ids, pos, w, c, q); + (void)Qwen3_5DenseModel::ForwardDense(ids, pos, w, c, q); + CHECK(w.lm_head_fp4.d_dequant_b.get() == first); + + // ...and the tower ran the same fallback GEMM without acquiring any. + for (const Qwen3_5DenseLayerWeights& lw : w.layers) { + CHECK_FALSE(lw.mlp.gate_proj_fp4.keep_dequant_b); + CHECK(lw.mlp.gate_proj_fp4.d_dequant_b == nullptr); + CHECK(lw.mlp.up_proj_fp4.d_dequant_b == nullptr); + CHECK(lw.mlp.down_proj_fp4.d_dequant_b == nullptr); + } + + // A BF16 head acquires none — the hook is inert off the packed path. + Qwen3_5DenseWeights bf16_w = MakeWeights(c); + bf16_w.lm_head = ReferenceBf16Head(f, N, K); + const std::unique_ptr bf16_model = + vllm::BorrowQwen3_5DenseLoadedModel(bf16_w); + ModelRegistry::Prepare(*bf16_model, c, q); + CHECK(bf16_w.lm_head_fp4.d_dequant_b == nullptr); +} + +// ── 8. The opt-in belongs to ONE LOADER, and the LOADER is what proves it ──── +// The case above builds its tower with `MakeNvfp4Weight` — direct struct +// construction — so it pins what the FORWARD does with a weight that did not opt +// in, and says nothing about which loader may set the flag. Adding +// `r.keep_dequant_b = true;` to `LoadNvfp4AnyNaming`, the one function every +// dense NVFP4 TOWER projection flows through (MLP gate/up/down via +// `LoadDenseMlp`, attention q/k/v/o via `LoadAttnDense`, GDN `out_proj` via +// `LoadGdnDense`), reopens the whole-tower bf16 expansion of #203 verbatim while +// every case above stays green — and the CUDA gate is blind by construction, +// because `kMatmulNvfp4` is registered there so `ResidentNvfp4DequantB` is never +// reached. +// +// So load a fully-NVFP4 layer of BOTH types, under BOTH namings, through the REAL +// loader, and sweep every `Nvfp4Weight` the layer struct owns. The sweep +// enumerates FIELDS, not call sites, so a fourth opt-in site anywhere in the +// dense loader fails it. +TEST_CASE("qwen27 dense lm_head: LoadDenseLmHead is the ONLY loader that opts a weight in") { + const HfConfig c = MakeConfig(); + const int64_t N = c.vocab_size, K = c.hidden_size; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 53); + + for (const bool ct_naming : {false, true}) { + CAPTURE(ct_naming); + Bag bag; // outlives every weight below: the bf16 arms BORROW its bytes + if (ct_naming) { + PutCtNvfp4Head(bag, "lm_head", N, K, f); + } else { + PutModelOptHead(bag, "lm_head", N, K, f, /*with_input_scale=*/true); + } + for (int64_t l = 0; l < c.num_hidden_layers; ++l) + PutNvfp4DenseLayer(bag, c, c.layer_types[static_cast(l)], l, + ct_naming); + + // The head, through the loader that IS allowed to opt in. + OwnedTensor unused_bf16; + Nvfp4Weight head; + LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", unused_bf16, head); + REQUIRE_FALSE(head.Empty()); + CHECK(head.keep_dequant_b); + + // The tower, through the loaders that are NOT. + int populated = 0; + for (int64_t l = 0; l < c.num_hidden_layers; ++l) { + const std::string type = c.layer_types[static_cast(l)]; + const Qwen3_5DenseLayerWeights lw = + vllm::LoadQwen3_5DenseLayer(bag.Resolver(), bag.Has(), type, l); + for (const auto& [name, w] : LayerNvfp4Weights(lw)) { + // An empty slot would pass the opt-in check for free, so count what the + // loader actually routed and require the census below: a fixture that + // stopped producing NVFP4 tower weights must not read as a pass. + if (w->Empty()) continue; + ++populated; + INFO(type << " layer " << l << " " << name); + CHECK_FALSE(w->keep_dequant_b); + } + } + // linear_attention: gdn.out_proj + mlp gate/up/down = 4. + // full_attention: attn q/k/v/o + mlp gate/up/down = 7. + CHECK(populated == 11); + } +} diff --git a/tests/vllm/models/test_qwen27_paged_forward.cpp b/tests/vllm/models/test_qwen27_paged_forward.cpp index 82654bd57..4b4559fdf 100644 --- a/tests/vllm/models/test_qwen27_paged_forward.cpp +++ b/tests/vllm/models/test_qwen27_paged_forward.cpp @@ -1341,3 +1341,67 @@ TEST_CASE("qwen27 dense paged: one-shot prefill == chunked prefill (state contin CHECK(d < 1e-4); } } + +// PERF-27B-LMHEAD-FP4 (issue #213). The PAGED forward has TWO lm_head call sites +// — gathered (prefill/mixed) and the non-gathered full [T,vocab] arm — and a +// PACKED head leaves the bf16 owner EMPTY, so reverting either hands the logits +// GEMM an empty tensor. test_qwen27_dense_lmhead_fp4 pins the NUMERICS but runs +// only the EAGER forward; this pins that both PAGED arms SELECT the packed head. +namespace { +Nvfp4Weight MakePackedHead(int64_t n, int64_t k, uint64_t seed) { + Nvfp4Weight w; + w.n = n; + w.k = k; + w.scale2 = 0.125F; // ModelOpt weight_scale_2 IS the scale + w.packed = MakeOwned(DType::kI8, {n, k / 2}, seed); + w.scale = MakeOwned(DType::kI8, {n, k / 16}, seed + 1); + // fp4 operands are RAW bytes, and MakeOwned has no kI8 arm — its f32 branch + // over-allocates 4 B/element — so size and fill them exactly here. + w.packed.bytes.resize(static_cast(n) * static_cast(k / 2)); + w.scale.bytes.resize(static_cast(n) * static_cast(k / 16)); + auto* pb = w.packed.bytes.data(); + for (size_t i = 0; i < w.packed.bytes.size(); ++i) + pb[i] = static_cast((i * 37U + 11U) & 0x77U); + const uint8_t kE4M3PowersOfTwo[4] = {0x34, 0x38, 0x3C, 0x40}; // .25 .5 1 2 + auto* sb = w.scale.bytes.data(); + for (size_t i = 0; i < w.scale.bytes.size(); ++i) + sb[i] = kE4M3PowersOfTwo[i & 3U]; + return w; +} +} // namespace + +TEST_CASE("qwen27 dense paged: both lm_head arms run a PACKED NVFP4 head") { + const HfConfig c = MakeConfig(); + const int64_t T = 5, V = c.vocab_size; + Qwen3_5DenseWeights w = MakeWeights(c); + w.lm_head_fp4 = MakePackedHead(V, c.hidden_size, 4242); + + const std::vector ids{3, 11, 7, 20, 5}, pos{0, 1, 2, 3, 4}; + vt::Queue q = Q(); + const std::vector eager = + Qwen3_5DenseModel::ForwardDense(ids, pos, w, c, q); + REQUIRE(eager.size() == static_cast(T * V)); + + const CommonAttentionMetadata am = PrefillAttnMeta(T, {0, 1}, 8, 0); + const GDNAttentionMetadata gm = PrefillGdnMeta(T, 0); + + // Non-gathered arm: empty logits_indices -> the full [T, vocab]. + { + CachePool pool(c, 4, 8); + const std::vector full = Qwen3_5DenseModel::Forward( + ids, pos, am, gm, pool.attn_kv, pool.gdn_state, w, c, q, {}); + REQUIRE(full.size() == eager.size()); + CHECK(MaxAbsDiff(full, eager, full.size()) < 1e-3); + } + // Gathered arm: the prefill shape the engine actually runs -> [1, vocab]. + { + CachePool pool(c, 4, 8); + const std::vector got = + Qwen3_5DenseModel::Forward(ids, pos, am, gm, pool.attn_kv, + pool.gdn_state, w, c, q, + {static_cast(T - 1)}); + REQUIRE(got.size() == static_cast(V)); + const std::vector tail(eager.end() - V, eager.end()); + CHECK(MaxAbsDiff(got, tail, got.size()) < 1e-3); + } +} diff --git a/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp b/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp index 076f5da0f..f556747bd 100644 --- a/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp +++ b/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp @@ -279,7 +279,12 @@ TEST_CASE("test_mtp_load_model_unified: dense MTP shares target embedding and lm CHECK_FALSE(model.has_own_lm_head()); CHECK(&model.embed_tokens() == &target.embed_tokens); CHECK(model.lm_head() == &target.lm_head); - CHECK(model.lm_head_fp4() == nullptr); + // PERF-27B-LMHEAD-FP4 (issue #213): the dense drafter shares the target's + // PACKED head too, so the pointer is the target's field rather than null. This + // target is bf16, so the field is EMPTY and the bf16 arm is still selected. + CHECK(model.lm_head_fp4() == &target.lm_head_fp4); + REQUIRE(model.lm_head_fp4() != nullptr); + CHECK(model.lm_head_fp4()->Empty()); CHECK(weights.NumLayers() == 1); REQUIRE(weights.dense_layers.size() == 1); CHECK(weights.fc.nk);