From e3a10952c71048edc98ceddf303a20af160f01fc Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 9 Aug 2026 22:09:00 +0000 Subject: [PATCH 1/5] spec(PERF-27B-LMHEAD-FP4): keep the ModelOpt NVFP4 lm_head packed The dense loader's U8 branch dequantizes a ModelOpt NVFP4 lm_head into a BF16 [in,out] operand, so the logits GEMM re-reads ~2.543 GB every decode step where the packed head is ~0.715 GB. The transposed storage additionally forces a row-major NN GEMM with no nvjet_sm121 kernel, which is why an SM80 CUTLASS tile is selected on an sm_121a part. vLLM keeps the head quantized: ModelOptMixedPrecisionConfig.get_quant_method accepts ParallelLMHead and _quantized_layer_prefix_candidates appends the bare lm_head key, then ModelOptNvFp4W4A16LinearMethod pins MarlinNvFp4LinearKernel. Spec only; no implementation. Records that test_qwen27_paged_engine cannot see this path because its checkpoint ships a BF16 head. Refs #213 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode] --- .agents/specs/perf-27b-lmhead-nvfp4.md | 142 +++++++++++++++++++++++++ 1 file changed, 142 insertions(+) create mode 100644 .agents/specs/perf-27b-lmhead-nvfp4.md diff --git a/.agents/specs/perf-27b-lmhead-nvfp4.md b/.agents/specs/perf-27b-lmhead-nvfp4.md new file mode 100644 index 000000000..22deb16a4 --- /dev/null +++ b/.agents/specs/perf-27b-lmhead-nvfp4.md @@ -0,0 +1,142 @@ +# 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` + +## 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:1359-1362` — vLLM **deletes** `input_scale` on the W4A16 path. +- `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`. Retire the stale "lm_head is never quantized" rule at `:506`. +3. `DenseLmHead` (`qwen3_5.cpp:1248-1250`) gains the packed branch so the tied + and `nk` cases keep one code path. Both consumers must route through it: + `:7024-7026` (gather) and `:7028-7030` (non-gather), plus the eager + `ForwardLogits` arm at `:6671-6674`. +4. Build the Marlin resident **pre-capture**, mirroring `:6424-6425`. A resident + built inside capture bakes a stack address and fails on replay. +5. Stop staging the BF16 head owner in `PrepareBf16Resident` (`:6456`), or the + ~2.3 GiB RSS win does not materialize. +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:1359-1362`). Keep `VT_MODELOPT_W4A4=0` and assert + `lm_head_fp4.IsTrueW4A4() == false`. +- **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. + +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; expect ~2.3 GiB lower with the head packed. + +## 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. + +## 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. From 5e515be901facf946e783bcf593cdbb619ab0851 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 9 Aug 2026 22:28:12 +0000 Subject: [PATCH 2/5] perf(PERF-27B-LMHEAD-FP4): keep the ModelOpt NVFP4 lm_head packed MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The dense loader's U8 branch dequantized a ModelOpt NVFP4 lm_head into a BF16 [in,out] operand, so the logits GEMM re-read ~2.543 GB every decode step where the packed head is ~0.715 GB, and the transposed [K,N] storage forced a row-major NN GEMM with no nvjet_sm121 kernel (an SM80 CUTLASS tile on an sm_121a part). vLLM keeps that head quantized: ModelOptMixedPrecisionConfig.get_quant_method accepts ParallelLMHead (modelopt.py:2508-2536) and _quantized_layer_prefix_candidates appends the bare `lm_head` key (modelopt.py:2491-2496), so ModelOptNvFp4W4A16LinearMethod — which pins MarlinNvFp4LinearKernel (modelopt.py:1249,1283-1284) and DELETES input_scale (modelopt.py:1359-1362) — resolves it, and logits_processor._apply_head (logits_processor.py:98-133) calls quant_method.apply every step. Verified against the pinned oracle at 555967922. LoadDenseLmHead routes an NVFP4 head through the SAME LoadNvfp4AnyNaming every other NVFP4 projection takes, into a new Qwen3_5DenseWeights::lm_head_fp4; the three dense consumers (eager ForwardDense, the gathered and non-gathered paged arms) and the dense MTP sibling now all select through one DenseLogitsF32D helper. The Marlin resident is built PRE-CAPTURE from the registry prepare hook (it Copies a host stack float; captured, that bakes a dangling stack address), and PrepareBf16Resident no longer stages a BF16 head owner. BF16, FP8, GGUF and tied heads are byte-unchanged, so every recorded unsloth benchmark is unaffected. VT_LMHEAD_FP4=0 is the same-binary rollback. test_qwen27_paged_engine cannot see this path — its checkpoint ships a BF16 head — hence the new synthetic loader+numerics gate. Refs #213 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode] --- docs/ENVIRONMENT.md | 1 + docs/FEATURES.md | 1 + docs/USAGE.md | 11 + .../model_executor/models/qwen3_5_dense.h | 41 +- src/vllm/model_executor/models/qwen3_5.cpp | 83 +++- .../model_executor/models/qwen3_5_dense.cpp | 7 +- .../models/qwen3_5_dense_weights.cpp | 41 +- tests/CMakeLists.txt | 5 + tests/parity/test_qwen27_dense_lmhead_fp4.cpp | 453 ++++++++++++++++++ .../v1/spec_decode/test_mtp_speculator.cpp | 8 +- 10 files changed, 624 insertions(+), 27 deletions(-) create mode 100644 tests/parity/test_qwen27_dense_lmhead_fp4.cpp diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 0c58b01bf..9ceed44dd 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`). `=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 c391d6774..737716f8d 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 gate PENDING | ✅ | ☐ | ☐ | | 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/USAGE.md b/docs/USAGE.md index ab22f9a03..4b8468830 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -271,6 +271,17 @@ 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`. 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..532fcf362 100644 --- a/include/vllm/model_executor/models/qwen3_5_dense.h +++ b/include/vllm/model_executor/models/qwen3_5_dense.h @@ -97,14 +97,25 @@ 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). A BF16 or FP8 head +// is materialized into `lm_head`; a ModelOpt/CT NVFP4 head 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 the MoE arm's Qwen3_5MoeWeights::lm_head_fp4 and + // vLLM's own decision to leave the head quantized: + // ModelOptMixedPrecisionConfig.get_quant_method accepts ParallelLMHead + // (modelopt.py:2508-2536) and _quantized_layer_prefix_candidates appends the + // bare `lm_head` key (modelopt.py:2491-2496), so ModelOptNvFp4W4A16LinearMethod + // — which pins 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 +146,21 @@ 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` (leaving `bf16_out` empty), every other storage form is materialized +// bf16 [in, out] into `bf16_out` by LoadLmHeadAnyDtype (leaving `fp4_out` empty). +// `proj` is the module path WITHOUT the trailing ".weight" (i.e. "lm_head"). +// Exported for the loader gate. +void LoadDenseLmHead(const TensorResolver& get, + const std::function& has, + const std::string& proj, OwnedTensor& bf16_out, + Nvfp4Weight& fp4_out); + +// 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, @@ -188,6 +214,13 @@ class Qwen3_5DenseModel { static void PrepareBf16Resident(const Qwen3_5DenseWeights& weights, vt::Queue& queue); + // PERF-27B-LMHEAD-FP4 (issue #213). Build the packed NVFP4 `lm_head_fp4` + // Marlin W4A16 resident, PRE-CAPTURE. Inert when the head is not packed or + // Marlin is not the selected path. Called from the registry `prepare` hook, + // mirroring Qwen3_5Model::PrepareMarlinResident. + static void PrepareMarlinResident(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/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 78b8d9aa4..c509ac686 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -2568,6 +2568,28 @@ DBuf MatmulNvfp4F32D(Dev d, const Tensor& x, const Nvfp4Weight& 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 +// 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's +// logits_processor._apply_head -> lm_head.quant_method.apply +// (logits_processor.py:98-133) and the MoE arm above; every dense consumer +// (eager ForwardDense, the gathered and non-gathered paged arms) routes here so +// exactly one head layout is ever selected. +// +// The bf16 arm keeps BOTH of its existing shapes: the tied embed_tokens raw +// [V,H] owner (nk -> MatmulBf16LogitsF32D) and the transposed [H,V] owner. +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). @@ -6452,6 +6474,33 @@ void Qwen3_5Model::PrepareMarlinResident(const Qwen3_5MoeWeights& weights, #endif } +// PERF-27B-LMHEAD-FP4 (issue #213). Build the PACKED dense head's Marlin W4A16 +// resident once, at prepare time — which is strictly BEFORE any decode-graph +// capture. This is not an optimization: BuildMarlinDenseResident Allocs, +// launches the repack, and Copies a HOST STACK float (the processed global +// scale) to the device. Run lazily from inside a captured region that bakes a +// dangling stack address into the graph and every replay reads freed memory. +// Same arm as Qwen3_5Model::PrepareMarlinResident's lm_head build above. +void Qwen3_5DenseModel::PrepareMarlinResident(const Qwen3_5DenseWeights& weights, + vt::Queue& 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.Empty() || weights.lm_head_fp4.IsTrueW4A4() || + !MarlinMoeEnabled() || + !vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, queue.device.type)) { + return; + } + Dev d{vt::GetBackend(queue.device.type), queue}; + BuildMarlinDenseResident(d, weights.lm_head_fp4, + MarlinDenseResidentFor(&weights.lm_head_fp4)); + d.b.Synchronize(d.q); +#else + (void)weights; + (void)queue; +#endif +} + void Qwen3_5DenseModel::PrepareBf16Resident( const Qwen3_5DenseWeights& weights, vt::Queue& queue) { VT_CHECK(platforms::GetPlatform(queue.device.type).needs_weight_staging(), @@ -6470,7 +6519,10 @@ void Qwen3_5DenseModel::PrepareBf16Resident( raw(weights.embed_tokens); raw(weights.final_norm); - raw(DenseLmHead(weights)); + // PERF-27B-LMHEAD-FP4: a PACKED head has no bf16 owner to stage, and staging + // one would hand back the ~2.3 GiB keeping it packed just saved. The packed + // resident is built by PrepareMarlinResident instead. + if (weights.lm_head_fp4.Empty()) raw(DenseLmHead(weights)); for (const Qwen3_5DenseLayerWeights& layer : weights.layers) { raw(layer.input_layernorm); raw(layer.post_attention_layernorm); @@ -6685,10 +6737,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; @@ -6700,7 +6751,13 @@ 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 — otherwise ForwardLogits would fall through to + // an empty bf16 owner on a ModelOpt NVFP4 checkpoint. Empty on every + // BF16/FP8/tied dense target, where the bf16 arm is selected exactly as + // before. Mirrors the MoE ctor below. + lm_head_fp4_(&target.lm_head_fp4) { VT_CHECK(weights.kind == Qwen3_5MTPKind::kDense, "qwen3_5 MTP: dense target requires dense MTP weights"); } @@ -7030,21 +7087,19 @@ 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, so a PACKED NVFP4 head + // (PERF-27B-LMHEAD-FP4) and the bf16/tied owner select the same way here as in + // the eager forward. 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..2594cf73b 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 Marlin resident + // HERE, before the runner ever captures a decode graph. Inert on every + // BF16/FP8/GGUF/tied dense checkpoint. Mirrors PrepareQwen3_5Moe. + auto& qwen = static_cast(model); + Qwen3_5DenseModel::PrepareMarlinResident(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..9af9d8388 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp @@ -219,11 +219,12 @@ 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 (identical call, no +// dequant). The NVFP4 form no longer reaches this function at all: since +// PERF-27B-LMHEAD-FP4 (issue #213) LoadDenseLmHead routes it to the PACKED +// `Qwen3_5DenseWeights::lm_head_fp4` instead, which is what vLLM does. 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 @@ -499,6 +500,31 @@ 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 + // `weight_scale_2`-is-the-scale vs compressed-tensors + // `weight_global_scale`-is-the-divisor split is handled in exactly one place, + // and `input_scale` is left unconsumed (IsTrueW4A4() stays false) unless + // VT_MODELOPT_W4A4=1. vLLM makes the same two decisions: the mixed scheme is + // designed to resolve a quantized head (modelopt.py:2491-2496,2508-2536), and + // ModelOptNvFp4W4A16LinearMethod DELETES input_scale (modelopt.py:1359-1362). + if (DenseLmHeadFp4Enabled() && IsNvfp4Projection(has, proj)) { + fp4_out = LoadNvfp4AnyNaming(get, has, proj); + return; + } + bf16_out = LoadLmHeadAnyDtype(get, has, proj + ".weight"); +} + OwnedTensor LoadMergedBf16RawNK(const TensorResolver& get, const std::vector& names) { // Extracted to the shared dense_weight_loaders.h (SEAM GAP #3); this retains @@ -633,7 +659,7 @@ Qwen3_5DenseWeights LoadQwen3_5Dense(const std::vector& shards, // 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"); + 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 +678,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: the direct-device + // staging path this gates 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 ccbfaea3f..a600f201d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1238,6 +1238,11 @@ vllm_cpp_add_test(test_qwen36_gguf_nvfp4_compute # the paged dense engine reproduces the pip-vLLM oracle greedy continuation # token-for-token via the fp4-resident W4A4 GEMM. See qwen27b-w4a4-notes.md §5. vllm_cpp_add_test(test_qwen27_paged_engine parity/test_qwen27_paged_engine.cpp) +# `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) target_compile_definitions(test_qwen27_paged_engine PRIVATE PARITY_GOLDENS_DIR="${CMAKE_SOURCE_DIR}/tests/parity/goldens") target_include_directories(test_qwen27_paged_engine PRIVATE 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..9dc457588 --- /dev/null +++ b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp @@ -0,0 +1,453 @@ +// 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 that head quantized: `ModelOptMixedPrecisionConfig.get_quant_method` +// accepts `ParallelLMHead` (modelopt.py:2508-2536) and +// `_quantized_layer_prefix_candidates` appends the bare `lm_head` key +// (modelopt.py:2491-2496), so `ModelOptNvFp4W4A16LinearMethod` — which pins +// `MarlinNvFp4LinearKernel` (modelopt.py:1249,1283-1284) — resolves the head and +// `logits_processor._apply_head` (logits_processor.py:98-133) calls +// `lm_head.quant_method.apply` every step. Nothing materializes BF16. +// +// These cases pin the LOADER ROUTING and the NUMERICS 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 "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/qwen3_5_dense.h" +#include "vllm/transformers_utils/hf_config.h" +#include "vt/backend.h" +#include "vt/dtype.h" + +using vllm::DenseMlpWeights; +using vllm::HfConfig; +using vllm::LoadDenseLmHead; +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 head fixture. The nibbles and the 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 itself, not the +// compressed-tensors 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 explicitly refuses on this + // head (modelopt.py:1359-1362 DELETES input_scale on the W4A16 path). + 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)); + } +} + +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; +} + +// --- the small synthetic dense model (shape scaffold from +// tests/vllm/models/test_qwen27_dense_forward.cpp; hidden/vocab widened so the +// head's K and N are Marlin-shaped on a CUDA queue) --- + +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 exact head: +// dequant to f32, round to bf16, transpose. Built here from the fixture's own +// exact values, NOT by calling either production dequant — so the numerical case +// below compares the packed forward against an independent reference rather than +// against a shared helper. +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 whole point: NOTHING is materialized to bf16. The old U8 branch produced + // 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. `lm_head.input_scale` must NOT flip the head to W4A4 ────────────────── +TEST_CASE("qwen27 dense lm_head: input_scale does not select the W4A4 GEMM") { + constexpr int64_t N = 256, K = 128; + const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 11); + 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()); + // vLLM DELETES input_scale on the W4A16 head path (modelopt.py:1359-1362); + // consuming it here would route the head to the fp4-activation GEMM. + CHECK_FALSE(fp4.IsTrueW4A4()); + CHECK(fp4.alpha == 0.0F); +} + +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, built from the + // fixture's own exact values (see ReferenceBf16Head). + 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 of each already-exact dequantized value plus + // GEMM accumulation order. (Marlin W4A16 tolerance band.) + CHECK(max_abs / scale < 2e-2); +} diff --git a/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp b/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp index 076f5da0f..0a91aee9c 100644 --- a/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp +++ b/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp @@ -279,7 +279,13 @@ 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 now shares the target's + // PACKED head as well, so the pointer is the target's field rather than null. + // This target is plain bf16, so the field is EMPTY and ForwardLogits still + // selects the bf16 arm — the pre-#213 behavior for every bf16 checkpoint. + 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); From e66527c4fe51731924ea366b5dcf59ac7c8c4d7f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Mon, 10 Aug 2026 07:41:52 +0000 Subject: [PATCH 3/5] fix(PERF-27B-LMHEAD-FP4): review findings on the packed NVFP4 lm_head MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Six defects around the packed head, from the row's fresh review. The head itself is unchanged: it still loads packed, still runs the Marlin W4A16 logits GEMM, and the reviewer's hardware run stands (focused 5/5, test_qwen27_paged_engine 235/235, greedy continuations byte-identical packed vs dequant on nvidia/Qwen3.6-27B-NVFP4@0893e160, peak RSS 21.06 -> 19.36 GiB). 1. A compressed-tensors-named NVFP4 head silently became W4A4. LoadCtNvfp4Raw sets alpha from `input_global_scale` UNCONDITIONALLY — correct for a 27B TOWER projection, wrong for an output head — so a CT-named head came back with alpha=0.0078125 and IsTrueW4A4()==true, contradicting both the code's own comment and docs/USAGE.md. That head would take the fp4-activation GEMM vLLM refuses for it AND make the pre-capture Marlin build early-return. LoadDenseLmHead now drops the activation globals on BOTH spellings unless VT_MODELOPT_W4A4=1, mirroring ModelOptNvFp4W4A16LinearMethod, which deletes input_scale (modelopt.py:1365; the placeholder is registered at :1358 — verified against the pinned oracle 555967922). The trap was masked only by a second bug: the model loader probed `lm_head.weight`, so a pure-CT head was read as tie_word_embeddings. DenseCheckpointHasLmHead now accepts either naming, which is what USAGE.md already advertised. 2. Every backend without an fp4 GEMM dequantized the whole head on EVERY forward call. MatmulNvfp4{F32,Bf16}D's fallback built a fresh K*N bf16 operand per call; CPU registers only kMatmulNvfp4Fp4, and Vulkan/Metal register neither candidate. At the gate model that is ~2.54 GB of allocate-and-dequantize per step where the OLD code paid it once at load. Nvfp4Weight gains a `d_dequant_b` resident, built once exactly like `d_packed`, and the registry prepare hook builds it up front — so the same hook now covers both arms and is named PrepareLmHeadResident. 3. The paged arms had no coverage. The numerical case ran only the eager ForwardDense, so reverting either paged lm_head call site to the bf16 owner (an EMPTY OwnedTensor on a packed head) left the full CPU suite green, as did dropping the pre-capture build. Both mutations are now RED. 4. Unreachable dead code credited with the win. PrepareBf16Resident's `if (lm_head_fp4.Empty()) raw(DenseLmHead(weights))` could never take the false branch: its only caller is reached under IsPlainBf16Qwen3_5Dense, which this row made false whenever the head is packed, and `raw` already skips an empty tensor. Guard deleted. The RSS win comes from the LOADER no longer building the f32 + bf16 arrays, not from anything skipped at staging time. 5. "~2.3 GiB" was wrong. bf16 head 2,543,206,400 B = 2.368 GiB; packed K*N/2 + K*N/16 = 715,264,000 B = 0.666 GiB; delta 1.70 GiB, which is exactly the measured 21.06 -> 19.36 GiB. Corrected in the spec and ENVIRONMENT.md. The 0.715 GB / 2.543 GB byte figures were already right. 6. Records. #213 added to the roadmap issue table (the third of the three places that must agree). FEATURES.md no longer says "supported and gated" and "CUDA gate PENDING" in the same cell. BENCHMARKS.md records the A/B: RSS -1.70 GiB and byte-identical continuations SOLID; throughput 11.197/11.193 packed vs 9.418/10.163 dequant is INDICATIVE ONLY — packed wins every leg and the packed legs agree to 0.04%, but the dequant legs disagree by 7.9%, so the direction is established and the magnitude is NOT, and a binding grid is owed. The spec records the IsQwen27QuantizedLinear deviation (zero production callers, so the design's retirement would move no behavior), fixes the stale :506 anchor to :527, and fixes the modelopt.py:1359-1362 anchor to :1365. The new test also no longer splits test_qwen27_paged_engine's configuration block. Also corrected: the capture-safety comment overstated the failure mode. Building the Marlin resident inside a capture would abort the capture with an error, not silently bake a dangling address; the build still belongs before capture. A BENCHMARKS measurement owes STATUS.md and NOW.md under the doc-checkpoint rule, and both pages were at zero headroom. Neither budget is widened and no checker is touched: the 27B-NVFP4 STATUS cell pays for its new fact by dropping two asides docs/BENCHMARKS.md states authoritatively — the "(ModelOpt FP8 tower)" parenthetical and the "decode ~100% GPU-busy" diagnostic, both on its Roof row — leaving that page byte-for-byte the same length, so its shrink-only ratchet stays byte-tight and untouched; the NOW row is compacted in place. Every measured number and binding claim is kept verbatim. Refs #213 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/roadmap_v1.md | 1 + .agents/specs/perf-27b-lmhead-nvfp4.md | 68 ++++++++--- docs/BENCHMARKS.md | 13 +++ docs/ENVIRONMENT.md | 2 +- docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- docs/USAGE.md | 10 +- .../model_executor/models/qwen3_5_dense.h | 18 ++- .../model_executor/models/qwen3_5_weights.h | 6 + src/vllm/model_executor/models/qwen3_5.cpp | 90 ++++++++++----- .../model_executor/models/qwen3_5_dense.cpp | 7 +- .../models/qwen3_5_dense_weights.cpp | 45 ++++++-- tests/CMakeLists.txt | 9 +- tests/parity/test_qwen27_dense_lmhead_fp4.cpp | 107 ++++++++++++++++++ .../vllm/models/test_qwen27_paged_forward.cpp | 60 ++++++++++ 16 files changed, 371 insertions(+), 71 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index b41ad14d4..a6a77cee9 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 exact bytes** at `776c56f1`; 95/95; raw-row guard | 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 tower native, tokens MATCH; #213 head packed, RSS -1.70 GiB | #213 grid; MLP marlin | | f32-out GEMV audit | **CLAIM WRONG**: 35B runs 41 `CastF32`/step (3.1%), a GATE model | Fold into the 35B lever | | Invocation-parity prevention | CI guard + checklist landing | build-verify `kGemvHeuristicAlgos` on dgx | | MiniMax-H3 lane | **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 37673c71b..e4ecfb24c 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -42,6 +42,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#242](https://github.com/mudler/vllm.cpp/issues/242) | — | `docs/FEATURES.md` drift: arch counts say 30 (registry has 35), multimodal-over-HTTP marked ☐ though W1-W3 landed | bug | | [#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 | | [#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 index 22deb16a4..da5557046 100644 --- a/.agents/specs/perf-27b-lmhead-nvfp4.md +++ b/.agents/specs/perf-27b-lmhead-nvfp4.md @@ -3,7 +3,8 @@ 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` +Base: `origin/main` @`04069bd7`; the review-findings round is rebased onto +`origin/main` @`723d96a8`. ## Scope @@ -49,7 +50,9 @@ an Ampere tile on an `sm_121a` part. - `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:1359-1362` — vLLM **deletes** `input_scale` on the W4A16 path. +- `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` / @@ -62,15 +65,30 @@ an Ampere tile on an `sm_121a` part. 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`. Retire the stale "lm_head is never quantized" rule at `:506`. -3. `DenseLmHead` (`qwen3_5.cpp:1248-1250`) gains the packed branch so the tied - and `nk` cases keep one code path. Both consumers must route through it: - `:7024-7026` (gather) and `:7028-7030` (non-gather), plus the eager - `ForwardLogits` arm at `:6671-6674`. -4. Build the Marlin resident **pre-capture**, mirroring `:6424-6425`. A resident - built inside capture bakes a stack address and fails on replay. -5. Stop staging the BF16 head owner in `PrepareBf16Resident` (`:6456`), or the - ~2.3 GiB RSS win does not materialize. + `: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 @@ -85,8 +103,16 @@ so every recorded `unsloth` benchmark is unaffected. 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:1359-1362`). Keep `VT_MODELOPT_W4A4=0` and assert - `lm_head_fp4.IsTrueW4A4() == false`. + (`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 @@ -104,6 +130,18 @@ RED first, `tests/parity/test_qwen27_dense_lmhead_fp4.cpp`: 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. + Port anchor: `marlin_utils_fp4.py` tolerances and shapes as used by the existing NVFP4A16 op tests. @@ -121,7 +159,9 @@ NVFP4A16 op tests. 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; expect ~2.3 GiB lower with the head packed. +- 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. ## Evidence diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 00eec219f..b65c8606c 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -94,6 +94,19 @@ 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**, SOLID | +| 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 9ceed44dd..a97474baf 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -90,7 +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`). `=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_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). 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 737716f8d..ef3ea3f02 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -65,7 +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 gate PENDING | ✅ | ☐ | ☐ | +| NVFP4 `lm_head` kept packed (no dequant at load) | ✅ `VT_LMHEAD_FP4` default-ON, #213; CUDA-gated on `nvidia/Qwen3.6-27B-NVFP4`@`0893e160` (greedy continuations byte-identical packed vs dequant, `test_qwen27_paged_engine` 235/235, RSS -1.70 GiB) | ✅ | ☐ | ☐ | | 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 1e664c7a4..144b151a5 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` now PACKED (#213, RSS -1.70 GiB) | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; decode 0.98x c1/c4 @`491c2f1e` after the warp-shuffle router, 0.87x c2, 0.92x c8. Async batch-1 token-0 degeneration FIXED (`VT_ASYNC_DEVICE_MIRROR` ON) | Token-exact SYNC+ASYNC (RED→GREEN); c16 0.93x; `VT_ASYNC_EXECUTOR` Option A (H2D out of capture) GREEN+RED but A/B NEUTRAL → OFF; c16 residual is prefill glue | | 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 4b8468830..9bc0c6485 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -278,9 +278,13 @@ On a Qwen3.6 dense checkpoint whose `lm_head` is stored NVFP4 (ModelOpt `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`. 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. +2.543 GB on `nvidia/Qwen3.6-27B-NVFP4` (measured peak host RSS 21.06 to 19.36 +GiB, a 1.70 GiB saving). 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 diff --git a/include/vllm/model_executor/models/qwen3_5_dense.h b/include/vllm/model_executor/models/qwen3_5_dense.h index 532fcf362..47dde0930 100644 --- a/include/vllm/model_executor/models/qwen3_5_dense.h +++ b/include/vllm/model_executor/models/qwen3_5_dense.h @@ -157,6 +157,13 @@ void LoadDenseLmHead(const TensorResolver& get, 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`. Exported so the gate can pin a CT-named +// head as such rather than as a tied one. +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(); @@ -214,11 +221,12 @@ class Qwen3_5DenseModel { static void PrepareBf16Resident(const Qwen3_5DenseWeights& weights, vt::Queue& queue); - // PERF-27B-LMHEAD-FP4 (issue #213). Build the packed NVFP4 `lm_head_fp4` - // Marlin W4A16 resident, PRE-CAPTURE. Inert when the head is not packed or - // Marlin is not the selected path. Called from the registry `prepare` hook, - // mirroring Qwen3_5Model::PrepareMarlinResident. - static void PrepareMarlinResident(const Qwen3_5DenseWeights& weights, + // 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 (so the + // forward never dequantizes per call). Inert when the head is not packed. + // Called from the registry `prepare` hook, mirroring the MoE sibling. + static void PrepareLmHeadResident(const Qwen3_5DenseWeights& weights, vt::Queue& queue); // Batched PAGED dense forward — the 27B analogue of Qwen3_5Model::Forward. diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h index 0e3ef6333..feda75f4e 100644 --- a/include/vllm/model_executor/models/qwen3_5_weights.h +++ b/include/vllm/model_executor/models/qwen3_5_weights.h @@ -202,6 +202,12 @@ struct Nvfp4Weight { // path. Uploaded once from the persistent `alpha` member; the diagnostic host // scalar path leaves this null. mutable std::shared_ptr d_alpha; + // Lazily-populated DEQUANTIZED bf16 [K=in, N=out] Matmul-B operand for the + // backends with NO fp4 GEMM (CPU / Vulkan / Metal fall through to `vt::Matmul` + // on a dequantized copy). Built ONCE and kept for the model lifetime like + // `d_packed`: per call it would rewrite K*N bf16 a step (~2.54 GB for the 27B + // head). Never populated on CUDA, where Marlin / vt::MatmulNvfp4 read packed. + mutable std::shared_ptr d_dequant_b; }; // Device-resident per-tensor FP8 (W8A8) weight — the 35B attn q/k/v/o + GDN diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index c509ac686..efa88bc45 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -1179,6 +1179,24 @@ 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). This is the fallback +// every backend without an fp4 GEMM takes — CPU registers only kMatmulNvfp4Fp4, +// Vulkan/Metal neither kMatmulNvfp4 nor the Marlin grouped GEMM — so uncached it +// rebuilds K*N bf16 per call. PERF-27B-LMHEAD-FP4 moved the 27B head here, where +// the loader used to pay that dequant once and per call it is ~2.54 GB a step. +Tensor ResidentNvfp4DequantB(Dev d, const Nvfp4Weight& w) { + 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}); +} + // 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, @@ -2545,7 +2563,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 @@ -2561,9 +2579,7 @@ 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()); + vt::Matmul(d.q, dout.t(), x, ResidentNvfp4DequantB(d, w)); } return dout; } @@ -2594,7 +2610,7 @@ DBuf DenseLogitsF32D(Dev d, const Tensor& x, const Qwen3_5DenseWeights& weights) // 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 @@ -2607,9 +2623,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()); + vt::Matmul(d.q, dout.t(), x, ResidentNvfp4DequantB(d, w)); } return dout; } @@ -6474,31 +6488,44 @@ void Qwen3_5Model::PrepareMarlinResident(const Qwen3_5MoeWeights& weights, #endif } -// PERF-27B-LMHEAD-FP4 (issue #213). Build the PACKED dense head's Marlin W4A16 -// resident once, at prepare time — which is strictly BEFORE any decode-graph -// capture. This is not an optimization: BuildMarlinDenseResident Allocs, -// launches the repack, and Copies a HOST STACK float (the processed global -// scale) to the device. Run lazily from inside a captured region that bakes a -// dangling stack address into the graph and every replay reads freed memory. -// Same arm as Qwen3_5Model::PrepareMarlinResident's lm_head build above. -void Qwen3_5DenseModel::PrepareMarlinResident(const Qwen3_5DenseWeights& weights, +// 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 (the processed global scale) whose source is a +// function-local temporary. CUDA aborts such a capture with an error rather +// than baking it silently, but a graph is not where a weight gets built. Same +// arm as Qwen3_5Model::PrepareMarlinResident's lm_head build above. +// +// Backends with NO fp4 GEMM (CPU / Vulkan / Metal): build the dequantized bf16 +// [K,N] operand here instead, so it is paid once, never 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.Empty() || weights.lm_head_fp4.IsTrueW4A4() || - !MarlinMoeEnabled() || - !vt::OpRegistered(vt::OpId::kMoeGroupedGemmNvfp4Marlin, queue.device.type)) { + 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; } - Dev d{vt::GetBackend(queue.device.type), queue}; - BuildMarlinDenseResident(d, weights.lm_head_fp4, - MarlinDenseResidentFor(&weights.lm_head_fp4)); - d.b.Synchronize(d.q); -#else - (void)weights; - (void)queue; #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( @@ -6519,10 +6546,13 @@ void Qwen3_5DenseModel::PrepareBf16Resident( raw(weights.embed_tokens); raw(weights.final_norm); - // PERF-27B-LMHEAD-FP4: a PACKED head has no bf16 owner to stage, and staging - // one would hand back the ~2.3 GiB keeping it packed just saved. The packed - // resident is built by PrepareMarlinResident instead. - if (weights.lm_head_fp4.Empty()) raw(DenseLmHead(weights)); + // PERF-27B-LMHEAD-FP4: `raw` is already a no-op for a PACKED head, whose bf16 + // owner is empty by construction (LoadDenseLmHead fills exactly one of the + // two), and this function is only reached under IsPlainBf16Qwen3_5Dense, false + // whenever the head is packed. The 1.70 GiB saving is the LOADER never + // building the f32 + bf16 arrays, not anything skipped here; the packed head's + // resident is built by PrepareLmHeadResident. + raw(DenseLmHead(weights)); for (const Qwen3_5DenseLayerWeights& layer : weights.layers) { raw(layer.input_layernorm); raw(layer.post_attention_layernorm); diff --git a/src/vllm/model_executor/models/qwen3_5_dense.cpp b/src/vllm/model_executor/models/qwen3_5_dense.cpp index 2594cf73b..11bc1ac75 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense.cpp @@ -103,11 +103,12 @@ std::unique_ptr LoadQwen3_5DenseModel( void PrepareQwen3_5Dense(LoadedModel& model, const HfConfig& config, vt::Queue& queue) { (void)config; - // PERF-27B-LMHEAD-FP4 (issue #213): build the packed lm_head's Marlin resident - // HERE, before the runner ever captures a decode graph. Inert on every + // PERF-27B-LMHEAD-FP4 (issue #213): build the packed lm_head's resident HERE — + // on CUDA before the runner ever captures a decode graph, and on a backend + // with no fp4 GEMM before the first forward pays the dequant. Inert on every // BF16/FP8/GGUF/tied dense checkpoint. Mirrors PrepareQwen3_5Moe. auto& qwen = static_cast(model); - Qwen3_5DenseModel::PrepareMarlinResident(qwen.weights(), queue); + 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 9af9d8388..135eac7ac 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 @@ -378,8 +386,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; @@ -513,18 +521,39 @@ void LoadDenseLmHead(const TensorResolver& get, const TensorExists& has, // PERF-27B-LMHEAD-FP4 (issue #213). An NVFP4 head stays PACKED, through the // SAME LoadNvfp4AnyNaming every other NVFP4 projection takes — so the ModelOpt // `weight_scale_2`-is-the-scale vs compressed-tensors - // `weight_global_scale`-is-the-divisor split is handled in exactly one place, - // and `input_scale` is left unconsumed (IsTrueW4A4() stays false) unless - // VT_MODELOPT_W4A4=1. vLLM makes the same two decisions: the mixed scheme is - // designed to resolve a quantized head (modelopt.py:2491-2496,2508-2536), and - // ModelOptNvFp4W4A16LinearMethod DELETES input_scale (modelopt.py:1359-1362). + // `weight_global_scale`-is-the-divisor split is handled in exactly one place. + // vLLM makes the same decision: the mixed scheme is designed to resolve a + // quantized head (modelopt.py:2491-2496,2508-2536). 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 on IsTrueW4A4(), silently skipping the + // pre-capture Marlin build. So drop the activation globals on BOTH spellings + // unless the VT_MODELOPT_W4A4 opt-in governing the ModelOpt arm is set. + if (!ModelOptW4A4OptIn()) { + fp4_out.input_global_scale_inv = 0.0F; + fp4_out.alpha = 0.0F; + } 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 @@ -658,7 +687,7 @@ 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")) { + if (DenseCheckpointHasLmHead(has, "lm_head")) { LoadDenseLmHead(get, has, "lm_head", w.lm_head, w.lm_head_fp4); } else { w.tied_lm_head = true; diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index a600f201d..15a200c36 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1238,15 +1238,16 @@ vllm_cpp_add_test(test_qwen36_gguf_nvfp4_compute # the paged dense engine reproduces the pip-vLLM oracle greedy continuation # token-for-token via the fp4-resident W4A4 GEMM. See qwen27b-w4a4-notes.md §5. vllm_cpp_add_test(test_qwen27_paged_engine parity/test_qwen27_paged_engine.cpp) +target_compile_definitions(test_qwen27_paged_engine PRIVATE + PARITY_GOLDENS_DIR="${CMAKE_SOURCE_DIR}/tests/parity/goldens") +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) -target_compile_definitions(test_qwen27_paged_engine PRIVATE - PARITY_GOLDENS_DIR="${CMAKE_SOURCE_DIR}/tests/parity/goldens") -target_include_directories(test_qwen27_paged_engine PRIVATE - ${CMAKE_SOURCE_DIR}/tests/parity ${CMAKE_SOURCE_DIR}/src) # 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 + diff --git a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp index 9dc457588..7c553ffb3 100644 --- a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp +++ b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp @@ -24,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -31,14 +32,17 @@ #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::DenseMlpWeights; using vllm::HfConfig; using vllm::LoadDenseLmHead; +using vllm::ModelRegistry; using vllm::Nvfp4Weight; using vllm::OwnedTensor; using vllm::Qwen3_5DenseLayerWeights; @@ -167,6 +171,24 @@ void PutModelOptHead(Bag& bag, const std::string& proj, int64_t n, int64_t k, } } +// The SAME fp4 bytes under compressed-tensors names. CT spells the global scale +// `weight_global_scale` and stores it as a DIVISOR (so 1/scale), and it ships a +// per-tensor `input_global_scale` next to every quantized Linear — which 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)); +} + uint16_t F32ToBf16(float v) { uint32_t bits = 0; std::memcpy(&bits, &v, sizeof(bits)); @@ -358,6 +380,44 @@ TEST_CASE("qwen27 dense lm_head: input_scale does not select the W4A4 GEMM") { CHECK(fp4.alpha == 0.0F); } +// ── 3b. The SAME rule under compressed-tensors names ───────────────────────── +// `LoadCtNvfp4Raw` consumes `input_global_scale` UNCONDITIONALLY, because on a +// TOWER projection the 27B CT checkpoint really is W4A4. An output head is not a +// tower projection: vLLM resolves it through `ModelOptNvFp4W4A16LinearMethod`, +// which DELETES `input_scale` (modelopt.py:1365), and USAGE.md advertises the CT +// spelling as kept packed on the same W4A16 path as the ModelOpt spelling. A +// W4A4 head would take the fp4-activation GEMM AND silently 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. USAGE.md advertises the CT spelling as kept + // packed; that is only true if the model loader looks for it. + CHECK(DenseCheckpointHasLmHead(bag.Has(), "lm_head")); + Bag tied; // a genuinely tied checkpoint ships no head under either naming + 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)); @@ -451,3 +511,50 @@ TEST_CASE("qwen27 dense lm_head: packed-head logits match the dequant-then-GEMM // 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 ───────────── +// 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 the LOADER used to build exactly once. Two ways +// to lose that, neither visible to a numerical assertion: rebuild it inside the +// GEMM (2.54 GB per decode step at the real 248320x5120), or drop the +// PrepareLmHeadResident call so the forward builds it — on CUDA that same call +// is the PRE-CAPTURE Marlin build. Pinned here by pointer identity. +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; + 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); + + // 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); +} diff --git a/tests/vllm/models/test_qwen27_paged_forward.cpp b/tests/vllm/models/test_qwen27_paged_forward.cpp index 82654bd57..978fef74b 100644 --- a/tests/vllm/models/test_qwen27_paged_forward.cpp +++ b/tests/vllm/models/test_qwen27_paged_forward.cpp @@ -1341,3 +1341,63 @@ 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 — the gathered (prefill/mixed) and the non-gathered full [T,vocab] arm — +// and a PACKED NVFP4 head leaves the bf16 owner EMPTY, so reverting either site +// to it hands the logits GEMM an empty tensor. The packed head's gate +// (test_qwen27_dense_lmhead_fp4) pins its NUMERICS but runs only the EAGER +// ForwardDense; this pins that both PAGED arms SELECT it. +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); + // MakeOwned fills f32/bf16 patterns; fp4 operands are raw bytes. + auto* pb = reinterpret_cast(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 = reinterpret_cast(w.scale.bytes.data()); + for (size_t i = 0; i < w.scale.bytes.size(); ++i) + sb[i] = kE4M3PowersOfTwo[i & 3U]; + return w; +} + +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); + } +} From cf9a3ba7317a216bc0671ee2a6ade816203a993c Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Mon, 10 Aug 2026 13:21:04 +0000 Subject: [PATCH 4/5] fix(PERF-27B-LMHEAD-FP4): bound the dequant-B resident to the lm_head Round-2 review findings on the packed NVFP4 head. The head itself is unchanged: it still loads packed and still runs the Marlin W4A16 logits GEMM on CUDA. 1. BLOCKER. The round-1 fix traded a per-call temporary for a LIFETIME resident on every non-CUDA backend, for every NVFP4 projection -- not just the head. The caching sat inside MatmulNvfp4{F32,Bf16}D, which also serve DenseMlpBlock gate/up/down, attention 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 and OOM-rebooting a Spark. CUDA is unaffected, which is why the CUDA gate could not see it. Residency is now a property of the WEIGHT, not of the GEMM: Nvfp4Weight::keep_dequant_b, default OFF, set by LoadDenseLmHead and by nothing else. It is worth its bytes exactly where one operand is re-read whole every step and there is one of it -- the output head. MatmulNvfp4DequantB carries the branch, so a weight that did not opt in keeps the PER-CALL temporary it had before this row, byte for byte. Alternatives weighed and rejected, recorded in the spec: a per-BACKEND switch (the tower is the problem on every fallback backend, not on 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). Footprint on a no-fp4-GEMM backend, per NVFP4 weight, steady state: before this row packed; per-call K*N*2 temporary round 1 packed + K*N*2 FOR EVERY ONE now packed, and K*N*2 for the HEAD ALONE Coverage that catches it: the prepare/residency case now populates an NVFP4 TOWER as well as an NVFP4 head and asserts that after two forwards no tower projection 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. 2. The claim was overstated, and it deepened a seam divergence. The caching landed on qwen3_5.cpp's PRIVATE dispatcher while the shared-seam dense_nvfp4::MatmulNvfp4W4A16D still rebuilt K*N bf16 per call. That parallel dispatcher predates this row -- dense_nvfp4_gemm.h was EXTRACTED from qwen3_5.cpp's anonymous namespace and its SCOPE comment records that the true-W4A4 path stays private there, so the two also carry independent Dev, MakeTensor, ResidentNvfp4 and DequantNvfp4ToBLayout copies. Unifying them is a refactor this row does not do; the exception is now recorded explicitly in the spec. What this row does instead is put the opt-in on the SHARED data type (Nvfp4Weight) that both dispatchers read, defaulted OFF, so they cannot disagree about a weight and no weight reachable from MatmulNvfp4W4A16D opts in today. This change covers the dense lm_head, and says so. 3. Anchor drift, in the file this row added. The input_scale delete was cited as modelopt.py:1359-1362 at two sites in the test; at the pinned oracle 555967922 that range is a blank line, the def and its docstring. Corrected to :1365 (the register stays :1358), re-verified against the oracle checkout. 4. docs/USAGE.md advertised "K*N/2 + K*N/16 instead of 2*K*N" unqualified. It now states that a backend with no fp4 GEMM pays the packed bytes PLUS one bf16 operand, built at prepare, and that only the head is kept that way. 5. Rebased onto origin/main a0fa12c7 (40 commits, incl. ENG-LOAD-DIRECT-UPLOAD #150). Keyed records verified against the target-branch version: NOW.md, roadmap_v1.md, STATUS.md and BENCHMARKS.md differ from origin/main by exactly this row's scoped edit and nothing else. Because #150 rewrote LoadCtNvfp4Raw to BORROW mmap'd bytes, the 21.06 -> 19.36 GiB reading was taken under different RSS accounting; it is now marked OWED a re-measurement in BENCHMARKS, ENVIRONMENT, FEATURES, STATUS, NOW and the spec rather than carried as accepted. No hardware was available to this session to re-take it. Also fixed, raised as judgement calls: the test helper MakePackedHead had external linkage (it sat outside its file's anonymous namespace) and built its fp4 operands through MakeOwned's f32 branch, over-allocating 4 B/element; it is now internal and exactly sized. IsQwen27QuantizedLinear's "lm_head is never quantized" rule is deliberately left alone -- deleting it and its checked-in expectation costs review budget this PR does not have, and the function still has zero production callers, so it moves no behavior. It is recorded in the spec. CPU Release ctest -j 1: 366/367. The one failure, test_serve_low_tools (test_online_gate_client), is PRE-EXISTING on main from the just-landed BENCH-SERVER-BINARY-RENAME work: tools/, tests/tools/, benchmarks/ and scripts/ are byte-identical to origin/main in this tree, so nothing here can reach it. Focused: test_qwen27_dense_lmhead_fp4 6/6, test_qwen27_paged_forward 20/20, test_mtp_speculator 13/13. check-pr-size: product 900/900. NO CUDA GATE WAS RUN in this session -- dgx.casa was not reachable -- so the CUDA rerun and the RSS re-measurement are owed to the operator. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/specs/perf-27b-lmhead-nvfp4.md | 54 ++++++- docs/BENCHMARKS.md | 2 +- docs/ENVIRONMENT.md | 2 +- docs/FEATURES.md | 2 +- docs/STATUS.md | 2 +- docs/USAGE.md | 10 +- .../model_executor/models/qwen3_5_dense.h | 37 ++--- .../model_executor/models/qwen3_5_weights.h | 14 +- src/vllm/model_executor/models/qwen3_5.cpp | 85 ++++++----- .../model_executor/models/qwen3_5_dense.cpp | 5 +- .../models/qwen3_5_dense_weights.cpp | 42 +++--- tests/parity/test_qwen27_dense_lmhead_fp4.cpp | 142 +++++++++--------- .../vllm/models/test_qwen27_paged_forward.cpp | 22 +-- .../v1/spec_decode/test_mtp_speculator.cpp | 7 +- 15 files changed, 242 insertions(+), 186 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index a6a77cee9..27f2f3200 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 exact bytes** at `776c56f1`; 95/95; raw-row guard | 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.85x**: FP8 tower native, tokens MATCH; #213 head packed, RSS -1.70 GiB | #213 grid; MLP marlin | +| 27B NVFP4 @`0893e160` | **0.85x**: FP8 tower native, tokens MATCH; #213 head packed | #213 grid + RSS re-measure post-#150 | | f32-out GEMV audit | **CLAIM WRONG**: 35B runs 41 `CastF32`/step (3.1%), a GATE model | Fold into the 35B lever | | Invocation-parity prevention | CI guard + checklist landing | build-verify `kGemvHeuristicAlgos` on dgx | | MiniMax-H3 lane | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | diff --git a/.agents/specs/perf-27b-lmhead-nvfp4.md b/.agents/specs/perf-27b-lmhead-nvfp4.md index da5557046..979b38283 100644 --- a/.agents/specs/perf-27b-lmhead-nvfp4.md +++ b/.agents/specs/perf-27b-lmhead-nvfp4.md @@ -3,8 +3,9 @@ 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 review-findings round is rebased onto -`origin/main` @`723d96a8`. +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 @@ -142,6 +143,14 @@ Added in the review-findings round, each pinned by a mutation that turns it RED: 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. + Port anchor: `marlin_utils_fp4.py` tolerances and shapes as used by the existing NVFP4A16 op tests. @@ -161,7 +170,10 @@ NVFP4A16 op tests. 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. + `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. ## Evidence @@ -169,6 +181,42 @@ NVFP4A16 op tests. 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 read, and default it +OFF — so the shared seam's fallback and the private one cannot disagree about a +weight, and no weight reachable from `MatmulNvfp4W4A16D` opts in today. The PR +body and commit message claim only the dense `lm_head`, never "every fp4 +projection". + ## Stop conditions - Stop and report `NEEDS_DECISION` if the packed head cannot be made diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index b65c8606c..4d6e70fd1 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -98,7 +98,7 @@ the same metric at higher concurrency (c8 p99 ITL 0.86x, but 1.055x at c16 and | Axis | Packed (`VT_LMHEAD_FP4=1`) | Dequant (`=0`) | Result | |---|---:|---:|---| -| Peak host RSS | 19.36 GiB | 21.06 GiB | **-1.70 GiB**, SOLID | +| 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 | | 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** | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index a97474baf..21ad01a86 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -90,7 +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). 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_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 ef3ea3f02..31ab33f2d 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -65,7 +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` (greedy continuations byte-identical packed vs dequant, `test_qwen27_paged_engine` 235/235, RSS -1.70 GiB) | ✅ | ☐ | ☐ | +| 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 144b151a5..ab79097dc 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 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` **0.85x BEHIND**; its NVFP4 `lm_head` now PACKED (#213, RSS -1.70 GiB) | +| 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; decode 0.98x c1/c4 @`491c2f1e` after the warp-shuffle router, 0.87x c2, 0.92x c8. Async batch-1 token-0 degeneration FIXED (`VT_ASYNC_DEVICE_MIRROR` ON) | Token-exact SYNC+ASYNC (RED→GREEN); c16 0.93x; `VT_ASYNC_EXECUTOR` Option A (H2D out of capture) GREEN+RED but A/B NEUTRAL → OFF; c16 residual is prefill glue | | 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 9bc0c6485..479b410a4 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -279,7 +279,15 @@ On a Qwen3.6 dense checkpoint whose `lm_head` is stored NVFP4 (ModelOpt 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). The head runs W4A16 under both namings: the on-disk +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. 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 diff --git a/include/vllm/model_executor/models/qwen3_5_dense.h b/include/vllm/model_executor/models/qwen3_5_dense.h index 47dde0930..13c4ea3d4 100644 --- a/include/vllm/model_executor/models/qwen3_5_dense.h +++ b/include/vllm/model_executor/models/qwen3_5_dense.h @@ -99,22 +99,21 @@ struct Qwen3_5DenseLayerWeights { // 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). A BF16 or FP8 head -// is materialized into `lm_head`; a ModelOpt/CT NVFP4 head stays PACKED in -// `lm_head_fp4` (PERF-27B-LMHEAD-FP4, issue #213). Exactly one is populated. +// 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 the MoE arm's Qwen3_5MoeWeights::lm_head_fp4 and - // vLLM's own decision to leave the head quantized: - // ModelOptMixedPrecisionConfig.get_quant_method accepts ParallelLMHead - // (modelopt.py:2508-2536) and _quantized_layer_prefix_candidates appends the - // bare `lm_head` key (modelopt.py:2491-2496), so ModelOptNvFp4W4A16LinearMethod - // — which pins 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. + // 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. @@ -148,10 +147,8 @@ OwnedTensor LoadLmHeadAnyDtype(const TensorResolver& get, // 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` (leaving `bf16_out` empty), every other storage form is materialized -// bf16 [in, out] into `bf16_out` by LoadLmHeadAnyDtype (leaving `fp4_out` empty). -// `proj` is the module path WITHOUT the trailing ".weight" (i.e. "lm_head"). -// Exported for the loader gate. +// `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, @@ -159,8 +156,7 @@ void LoadDenseLmHead(const TensorResolver& get, // True when the checkpoint ships an EXPLICIT head under either naming // (`.weight`, or `.weight_packed` for compressed-tensors NVFP4); -// false means `tie_word_embeddings`. Exported so the gate can pin a CT-named -// head as such rather than as a tied one. +// false means `tie_word_embeddings`. bool DenseCheckpointHasLmHead(const std::function& has, const std::string& proj); @@ -223,9 +219,8 @@ class Qwen3_5DenseModel { // 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 (so the - // forward never dequantizes per call). Inert when the head is not packed. - // Called from the registry `prepare` hook, mirroring the MoE sibling. + // 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); diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h index feda75f4e..5f2a0ccb2 100644 --- a/include/vllm/model_executor/models/qwen3_5_weights.h +++ b/include/vllm/model_executor/models/qwen3_5_weights.h @@ -202,11 +202,15 @@ struct Nvfp4Weight { // path. Uploaded once from the persistent `alpha` member; the diagnostic host // scalar path leaves this null. mutable std::shared_ptr d_alpha; - // Lazily-populated DEQUANTIZED bf16 [K=in, N=out] Matmul-B operand for the - // backends with NO fp4 GEMM (CPU / Vulkan / Metal fall through to `vt::Matmul` - // on a dequantized copy). Built ONCE and kept for the model lifetime like - // `d_packed`: per call it would rewrite K*N bf16 a step (~2.54 GB for the 27B - // head). Never populated on CUDA, where Marlin / vt::MatmulNvfp4 read packed. + // 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; }; diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index efa88bc45..b88be8c68 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -1180,12 +1180,11 @@ std::vector DequantNvfp4ToBLayout(const Nvfp4Weight& w) { } // The SAME bf16 [K=in, N=out] operand, uploaded ONCE and kept resident on the -// weight (mirror of ResidentNvfp4, same Backend deleter). This is the fallback -// every backend without an fp4 GEMM takes — CPU registers only kMatmulNvfp4Fp4, -// Vulkan/Metal neither kMatmulNvfp4 nor the Marlin grouped GEMM — so uncached it -// rebuilds K*N bf16 per call. PERF-27B-LMHEAD-FP4 moved the 27B head here, where -// the loader used to pay that dequant once and per call it is ~2.54 GB a step. +// 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); @@ -1197,6 +1196,22 @@ Tensor ResidentNvfp4DequantB(Dev d, const Nvfp4Weight& w) { 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, @@ -2579,25 +2594,20 @@ 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 { - vt::Matmul(d.q, dout.t(), x, ResidentNvfp4DequantB(d, w)); + 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 -// 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's -// logits_processor._apply_head -> lm_head.quant_method.apply -// (logits_processor.py:98-133) and the MoE arm above; every dense consumer -// (eager ForwardDense, the gathered and non-gathered paged arms) routes here so -// exactly one head layout is ever selected. -// -// The bf16 arm keeps BOTH of its existing shapes: the tied embed_tokens raw -// [V,H] owner (nk -> MatmulBf16LogitsF32D) and the transposed [H,V] owner. +// 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); @@ -2623,7 +2633,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 { - vt::Matmul(d.q, dout.t(), x, ResidentNvfp4DequantB(d, w)); + MatmulNvfp4DequantB(d, dout.t(), x, w); } return dout; } @@ -6492,15 +6502,12 @@ void Qwen3_5Model::PrepareMarlinResident(const Qwen3_5MoeWeights& weights, // 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 (the processed global scale) whose source is a -// function-local temporary. CUDA aborts such a capture with an error rather -// than baking it silently, but a graph is not where a weight gets built. Same -// arm as Qwen3_5Model::PrepareMarlinResident's lm_head build above. -// -// Backends with NO fp4 GEMM (CPU / Vulkan / Metal): build the dequantized bf16 -// [K,N] operand here instead, so it is paid once, never on the forward path. +// 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; @@ -6546,12 +6553,9 @@ void Qwen3_5DenseModel::PrepareBf16Resident( raw(weights.embed_tokens); raw(weights.final_norm); - // PERF-27B-LMHEAD-FP4: `raw` is already a no-op for a PACKED head, whose bf16 - // owner is empty by construction (LoadDenseLmHead fills exactly one of the - // two), and this function is only reached under IsPlainBf16Qwen3_5Dense, false - // whenever the head is packed. The 1.70 GiB saving is the LOADER never - // building the f32 + bf16 arrays, not anything skipped here; the packed head's - // resident is built by PrepareLmHeadResident. + // 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); @@ -6782,11 +6786,8 @@ Qwen3_5MTPModel::Qwen3_5MTPModel(const Qwen3_5MTPWeights& weights, config_(&config), embed_tokens_(&target.embed_tokens), lm_head_(&DenseLmHead(target)), - // PERF-27B-LMHEAD-FP4: the drafter shares the TARGET's head, so it must - // see the packed one too — otherwise ForwardLogits would fall through to - // an empty bf16 owner on a ModelOpt NVFP4 checkpoint. Empty on every - // BF16/FP8/tied dense target, where the bf16 arm is selected exactly as - // before. Mirrors the MoE ctor below. + // 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"); @@ -7117,10 +7118,8 @@ static DBuf DenseForwardLayers(Dev d, const Tensor& hidden_in, } // Logits gather-before-lm_head (prefill/mixed): same semantics as the 35B path. - // Both arms route through DenseLogitsF32D, so a PACKED NVFP4 head - // (PERF-27B-LMHEAD-FP4) and the bf16/tied owner select the same way here as in - // the eager forward. 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) { diff --git a/src/vllm/model_executor/models/qwen3_5_dense.cpp b/src/vllm/model_executor/models/qwen3_5_dense.cpp index 11bc1ac75..d212738d8 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense.cpp @@ -104,9 +104,8 @@ void PrepareQwen3_5Dense(LoadedModel& model, const HfConfig& config, vt::Queue& queue) { (void)config; // PERF-27B-LMHEAD-FP4 (issue #213): build the packed lm_head's resident HERE — - // on CUDA before the runner ever captures a decode graph, and on a backend - // with no fp4 GEMM before the first forward pays the dequant. Inert on every - // BF16/FP8/GGUF/tied dense checkpoint. Mirrors PrepareQwen3_5Moe. + // 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); } 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 135eac7ac..dc904bf3d 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp @@ -228,11 +228,10 @@ Nvfp4Weight LoadCtNvfp4Raw(const TensorResolver& get, const std::string& proj) { // a ModelOpt NVFP4 head; both hit the old unconditional BF16 assert. // // 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 (identical call, no -// dequant). The NVFP4 form no longer reaches this function at all: since -// PERF-27B-LMHEAD-FP4 (issue #213) LoadDenseLmHead routes it to the PACKED -// `Qwen3_5DenseWeights::lm_head_fp4` instead, which is what vLLM does. The U8 -// branch below survives ONLY as the VT_LMHEAD_FP4=0 in-binary rollback. +// 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 @@ -518,30 +517,29 @@ void LoadDenseLmHead(const TensorResolver& get, const TensorExists& has, 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 - // `weight_scale_2`-is-the-scale vs compressed-tensors - // `weight_global_scale`-is-the-divisor split is handled in exactly one place. - // vLLM makes the same decision: the mixed scheme is designed to resolve a - // quantized head (modelopt.py:2491-2496,2508-2536). + // 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 on IsTrueW4A4(), silently skipping the - // pre-capture Marlin build. So drop the activation globals on BOTH spellings - // unless the VT_MODELOPT_W4A4 opt-in governing the ModelOpt arm is set. + // `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"); @@ -707,8 +705,8 @@ 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: the direct-device - // staging path this gates only knows how to stage OwnedTensors. + // 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() || diff --git a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp index 7c553ffb3..54ebdb8bf 100644 --- a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp +++ b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp @@ -1,20 +1,13 @@ // 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). // -// vLLM keeps that head quantized: `ModelOptMixedPrecisionConfig.get_quant_method` -// accepts `ParallelLMHead` (modelopt.py:2508-2536) and -// `_quantized_layer_prefix_candidates` appends the bare `lm_head` key -// (modelopt.py:2491-2496), so `ModelOptNvFp4W4A16LinearMethod` — which pins -// `MarlinNvFp4LinearKernel` (modelopt.py:1249,1283-1284) — resolves the head and -// `logits_processor._apply_head` (logits_processor.py:98-133) calls -// `lm_head.quant_method.apply` every step. Nothing materializes BF16. -// -// These cases pin the LOADER ROUTING and the NUMERICS of that decision. They are -// synthetic (no checkpoint, no GPU) on purpose: the 235/235 +// 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 @@ -39,7 +32,6 @@ #include "vt/dtype.h" using vllm::DenseCheckpointHasLmHead; -using vllm::DenseMlpWeights; using vllm::HfConfig; using vllm::LoadDenseLmHead; using vllm::ModelRegistry; @@ -100,14 +92,11 @@ float RandV(uint64_t seed) { return static_cast(u * 0.16 - 0.08); } -// One deterministic ModelOpt NVFP4 head fixture. The nibbles and the 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: -// +// 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 itself, not the -// compressed-tensors divisor. +// `weight_scale_2` is the ModelOpt convention: the SCALE, not the CT divisor. constexpr float kFixtureScale2 = 0.125F; struct Nvfp4Fixture { @@ -161,9 +150,8 @@ void PutModelOptHead(Bag& bag, const std::string& proj, int64_t n, int64_t k, 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 explicitly refuses on this - // head (modelopt.py:1359-1362 DELETES input_scale on the W4A16 path). + // 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); @@ -171,10 +159,9 @@ void PutModelOptHead(Bag& bag, const std::string& proj, int64_t n, int64_t k, } } -// The SAME fp4 bytes under compressed-tensors names. CT spells the global scale -// `weight_global_scale` and stores it as a DIVISOR (so 1/scale), and it ships a -// per-tensor `input_global_scale` next to every quantized Linear — which is -// exactly the activation divisor that must NOT be consumed on an output head. +// 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}); @@ -189,6 +176,23 @@ void PutCtNvfp4Head(Bag& bag, const std::string& proj, int64_t n, int64_t k, 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; +} + uint16_t F32ToBf16(float v) { uint32_t bits = 0; std::memcpy(&bits, &v, sizeof(bits)); @@ -216,8 +220,7 @@ Fake MakeBf16(const std::vector& shape, uint64_t seed) { } // --- the small synthetic dense model (shape scaffold from -// tests/vllm/models/test_qwen27_dense_forward.cpp; hidden/vocab widened so the -// head's K and N are Marlin-shaped on a CUDA queue) --- +// test_qwen27_dense_forward.cpp; hidden/vocab widened to Marlin-shaped) --- OwnedTensor MakeOwned(DType dt, std::vector shape, uint64_t seed) { OwnedTensor t; @@ -316,11 +319,10 @@ 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 exact head: -// dequant to f32, round to bf16, transpose. Built here from the fixture's own -// exact values, NOT by calling either production dequant — so the numerical case -// below compares the packed forward against an independent reference rather than -// against a shared helper. +// 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; @@ -351,8 +353,8 @@ TEST_CASE("qwen27 dense lm_head: a ModelOpt NVFP4 head stays PACKED (no bf16 own LoadDenseLmHead(bag.Resolver(), bag.Has(), "lm_head", bf16, fp4); REQUIRE_FALSE(fp4.Empty()); - // The whole point: NOTHING is materialized to bf16. The old U8 branch produced - // a [K,N] bf16 owner (~2.543 GB at the real 248320x5120). + // 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); @@ -361,33 +363,18 @@ TEST_CASE("qwen27 dense lm_head: a ModelOpt NVFP4 head stays PACKED (no bf16 own 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. `lm_head.input_scale` must NOT flip the head to W4A4 ────────────────── -TEST_CASE("qwen27 dense lm_head: input_scale does not select the W4A4 GEMM") { - constexpr int64_t N = 256, K = 128; - const Nvfp4Fixture f = MakeNvfp4Fixture(N, K, 11); - 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()); - // vLLM DELETES input_scale on the W4A16 head path (modelopt.py:1359-1362); - // consuming it here would route the head to the fp4-activation GEMM. + // ── 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, because on a -// TOWER projection the 27B CT checkpoint really is W4A4. An output head is not a -// tower projection: vLLM resolves it through `ModelOptNvFp4W4A16LinearMethod`, -// which DELETES `input_scale` (modelopt.py:1365), and USAGE.md advertises the CT -// spelling as kept packed on the same W4A16 path as the ModelOpt spelling. A -// W4A4 head would take the fp4-activation GEMM AND silently skip the pre-capture -// Marlin build (which early-returns on IsTrueW4A4()). +// `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); @@ -396,10 +383,9 @@ TEST_CASE("qwen27 dense lm_head: a compressed-tensors NVFP4 head is W4A16, not W // 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. USAGE.md advertises the CT spelling as kept - // packed; that is only true if the model loader looks for it. + // logits off the embedding table. CHECK(DenseCheckpointHasLmHead(bag.Has(), "lm_head")); - Bag tied; // a genuinely tied checkpoint ships no head under either naming + 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")); @@ -484,8 +470,7 @@ TEST_CASE("qwen27 dense lm_head: packed-head logits match the dequant-then-GEMM const std::vector pos{0, 1, 2, 3}; vt::Queue q = CpuQ(); - // Reference arm: the bf16 [K,N] owner the OLD loader produced, built from the - // fixture's own exact values (see ReferenceBf16Head). + // 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 = @@ -507,19 +492,21 @@ TEST_CASE("qwen27 dense lm_head: packed-head logits match the dequant-then-GEMM 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 of each already-exact dequantized value plus - // GEMM accumulation order. (Marlin W4A16 tolerance band.) + // 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 ───────────── +// ── 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 the LOADER used to build exactly once. Two ways -// to lose that, neither visible to a numerical assertion: rebuild it inside the -// GEMM (2.54 GB per decode step at the real 248320x5120), or drop the -// PrepareLmHeadResident call so the forward builds it — on CUDA that same call -// is the PRE-CAPTURE Marlin build. Pinned here by pointer identity. +// 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; @@ -534,6 +521,13 @@ TEST_CASE("qwen27 dense lm_head: the packed head's resident is built once, at pr 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. @@ -550,6 +544,14 @@ TEST_CASE("qwen27 dense lm_head: the packed head's resident is built once, at pr (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); diff --git a/tests/vllm/models/test_qwen27_paged_forward.cpp b/tests/vllm/models/test_qwen27_paged_forward.cpp index 978fef74b..4b4559fdf 100644 --- a/tests/vllm/models/test_qwen27_paged_forward.cpp +++ b/tests/vllm/models/test_qwen27_paged_forward.cpp @@ -1342,12 +1342,12 @@ TEST_CASE("qwen27 dense paged: one-shot prefill == chunked prefill (state contin } } -// PERF-27B-LMHEAD-FP4 (issue #213). The PAGED forward has TWO lm_head call -// sites — the gathered (prefill/mixed) and the non-gathered full [T,vocab] arm — -// and a PACKED NVFP4 head leaves the bf16 owner EMPTY, so reverting either site -// to it hands the logits GEMM an empty tensor. The packed head's gate -// (test_qwen27_dense_lmhead_fp4) pins its NUMERICS but runs only the EAGER -// ForwardDense; this pins that both PAGED arms SELECT it. +// 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; @@ -1355,16 +1355,20 @@ Nvfp4Weight MakePackedHead(int64_t n, int64_t k, uint64_t seed) { 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); - // MakeOwned fills f32/bf16 patterns; fp4 operands are raw bytes. - auto* pb = reinterpret_cast(w.packed.bytes.data()); + // 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 = reinterpret_cast(w.scale.bytes.data()); + 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(); diff --git a/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp b/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp index 0a91aee9c..f556747bd 100644 --- a/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp +++ b/tests/vllm/v1/spec_decode/test_mtp_speculator.cpp @@ -279,10 +279,9 @@ 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); - // PERF-27B-LMHEAD-FP4 (issue #213): the dense drafter now shares the target's - // PACKED head as well, so the pointer is the target's field rather than null. - // This target is plain bf16, so the field is EMPTY and ForwardLogits still - // selects the bf16 arm — the pre-#213 behavior for every bf16 checkpoint. + // 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()); From fee4d66440dcdc1551ec55d1ff318d0a9c5db912 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Mon, 10 Aug 2026 16:43:23 +0000 Subject: [PATCH 5/5] test(PERF-27B-LMHEAD-FP4): pin the residency opt-in to ONE loader (#213) Round-3 review finding, test-only on the product side. The head is unchanged: it still loads packed and still runs the Marlin W4A16 logits GEMM on CUDA. The only non-test edit is a header DECLARATION of a function that already existed. 1. FINDING. `Nvfp4Weight::keep_dequant_b` is documented as "set by LoadDenseLmHead and by nothing else", and the round-2 fix is correct as written -- but nothing enforced the scope. The reviewer added ONE line to `LoadNvfp4AnyNaming` (qwen3_5_dense_weights.cpp:377): Nvfp4Weight r; r.keep_dequant_b = true; // a second setter and every gate stayed green: test_qwen27_dense_lmhead_fp4 6/6 1170/1170, test_qwen27_paged_forward 20/20 702/702, test_mtp_speculator 13/13 171/171. That function is what 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 the mutation reopens the round-2 blocker verbatim: a lifetime bf16 expansion of the whole tower on CPU, Vulkan, Metal, HIP and Tenstorrent. It escaped for two reasons: the residency case built its tower with MakeNvfp4Weight (direct struct construction, never a loader), and NO test in the repo called LoadDenseMlp, LoadDenseAttn or LoadQwen3_5DenseWeights at all. The CUDA gate is blind by construction -- kMatmulNvfp4 is registered on CUDA, so ResidentNvfp4DequantB is never reached there. This was the THIRD instance of one pattern on this row: the code correct at every site it is tested at, its SCOPE unpinned. So the new case pins the INVARIANT, not the call site. It loads a fully-NVFP4 dense layer of BOTH layer types, under BOTH namings, through the real `LoadQwen3_5DenseLayer`, and sweeps every Nvfp4Weight the layer STRUCT owns -- fields, not call sites -- asserting the flag is false on all of them and true on the head. A census assertion (11 routed projections per naming) keeps a fixture that stopped producing NVFP4 weights from reading as a pass. RED, three mutations, each restored byte-for-byte afterwards: M1 keep_dequant_b in LoadNvfp4AnyNaming's ModelOpt arm 11 failed M2 keep_dequant_b in LoadCtNvfp4Raw 11 failed M3 a FOURTH setter, one line after LoadDenseMlp's down_proj load -- outside LoadNvfp4AnyNaming 4 failed GREEN restored: 7/7, 1314/1314. `LoadQwen3_5DenseLayer` gains a `has`-taking overload in the header. The resolver-only overload answers `has` with a constant true, which forces every projection down the compressed-tensors spelling and cannot reach the ModelOpt arm at all. The definition already existed; only the declaration is new. 2. OVERSTATED CLAIM. The spec said the shared `Nvfp4Weight` opt-in means "the two dispatchers cannot disagree about a weight". They can: dense_nvfp4_gemm.h:626- 631 ignores the flag outright and rebuilds the per-call temporary for every weight. They merely do not, because nothing reachable there sets it -- and the divergence direction is benign, the shared seam UNDER-caches and so cannot reintroduce the whole-tower expansion. Reworded to that. 3. THE CPU FOOTPRINT DIRECTION WAS UNDISCLOSED. Every record carried only the CUDA -1.70 GiB. On a no-fp4-GEMM backend the head holds 0.666 GiB packed PLUS a 2.368 GiB bf16 operand. Arithmetic from the same K*N, now stated in docs/USAGE.md, docs/BENCHMARKS.md and the spec: CUDA 2.368 -> 0.666 -1.70 GiB Vulkan 2.368 + 2.368 -> 0.666 + 2.368 -1.70 GiB CPU 2.368 -> 0.666 + 2.368 +0.67 GiB So #203's backend genuinely improves and plain CPU genuinely regresses by the packed head's own bytes, paid once instead of rebuilding 2.368 GiB per step. 4. A BACKEND-ASYMMETRIC HARD ABORT, now recorded in the spec. 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 (CUDA returns at :6533 first). Unreachable in production; tests/vllm/models/test_qwen27_paged_forward.cpp:1377 constructs exactly such a weight and passes only because it never calls Prepare. The -1.70 GiB CUDA figure stays marked OWED, not carried: #150 rewrote LoadCtNvfp4Raw to borrow mmap'd bytes and changed the RSS accounting. Gates: clean CPU Release build, 0 warnings 0 errors. Focused test_qwen27_dense_lmhead_fp4 7/7 1314/1314 (was 6/6 1170/1170), test_qwen27_paged_forward 20/20 702/702, test_mtp_speculator 13/13 171/171. Full `ctest -j 1`: 369/369, 0 failed (round 2's one failure, test_serve_low_tools, was pre-existing on main and is fixed by the merge). scripts/agent-preflight.sh: all gates green. NO CUDA GATE WAS RUN and none is owed by this change -- it touches no CUDA behavior; the operator's run on cf9a3ba7 (paged_engine 235/235, paged_forward 702/702, dense_forward 333/333) stands. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/specs/perf-27b-lmhead-nvfp4.md | 81 +++++++++- docs/BENCHMARKS.md | 1 + docs/STATUS.md | 2 +- docs/USAGE.md | 7 +- .../model_executor/models/qwen3_5_dense.h | 11 ++ tests/parity/test_qwen27_dense_lmhead_fp4.cpp | 144 ++++++++++++++++++ 7 files changed, 242 insertions(+), 6 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 7b70b8be2..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; #213 head packed | #213 grid + RSS re-measure post-#150 | +| 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/specs/perf-27b-lmhead-nvfp4.md b/.agents/specs/perf-27b-lmhead-nvfp4.md index 979b38283..8bd023368 100644 --- a/.agents/specs/perf-27b-lmhead-nvfp4.md +++ b/.agents/specs/perf-27b-lmhead-nvfp4.md @@ -151,6 +151,35 @@ Added in the ROUND-2 findings round: 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. @@ -175,6 +204,22 @@ NVFP4A16 op tests. 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, @@ -211,12 +256,42 @@ 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 read, and default it -OFF — so the shared seam's fallback and the private one cannot disagree about a -weight, and no weight reachable from `MatmulNvfp4W4A16D` opts in today. The PR +(`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 diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index baef83383..de451c689 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -99,6 +99,7 @@ the same metric at higher concurrency (c8 p99 ITL 0.86x, but 1.055x at c16 and | 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** | diff --git a/docs/STATUS.md b/docs/STATUS.md index a1e914f55..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; its NVFP4 `lm_head` is kept PACKED (#213; RSS re-measure due) | +| 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 1fce04da8..f4b4f0af3 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -299,7 +299,12 @@ GiB, a 1.70 GiB saving on CUDA; the figure is owed a re-measurement after 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. Only the head is kept that way; every +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 diff --git a/include/vllm/model_executor/models/qwen3_5_dense.h b/include/vllm/model_executor/models/qwen3_5_dense.h index 13c4ea3d4..6b6279623 100644 --- a/include/vllm/model_executor/models/qwen3_5_dense.h +++ b/include/vllm/model_executor/models/qwen3_5_dense.h @@ -190,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 diff --git a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp index 54ebdb8bf..8808f83f7 100644 --- a/tests/parity/test_qwen27_dense_lmhead_fp4.cpp +++ b/tests/parity/test_qwen27_dense_lmhead_fp4.cpp @@ -193,6 +193,23 @@ Nvfp4Weight MakeNvfp4Weight(int64_t n, int64_t k, uint64_t seed) { 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)); @@ -219,6 +236,71 @@ Fake MakeBf16(const std::vector& shape, uint64_t seed) { 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) --- @@ -560,3 +642,65 @@ TEST_CASE("qwen27 dense lm_head: the packed head's resident is built once, at pr 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); + } +}