perf(PERF-27B-LMHEAD-FP4): keep the ModelOpt NVFP4 lm_head packed (#213)

.agents/NOW.md

@@ -15,7 +15,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`.
 | `SPEC-DSPARK` | **WORKS on 35B**: ON==OFF 48/48. ★fixed engine-wide draft-drop | Draft step ~6x a target step |
 | State record (#166) | **157 imports = 3,231,342 bytes** at `776c56f1`; 95/95 | Force-update #166; rerun readiness |
 | Laguna NVFP4 / DeepSeek-V4 decode | **CLOSED, byte-exact, default-ON**: 1.03x vLLM, 1.144x ds4 | Laguna vLLM K-run |
-| 27B NVFP4 @`0893e160` | **0.72x -> 0.85x**: FP8 tower native, tokens MATCH, RSS -3.2 GiB | NVFP4 MLP marlin, 68% of roof |
+| 27B NVFP4 @`0893e160` | **0.85x**: FP8 native, tokens MATCH; #213 head packed | #213 grid + RSS remeasure post-#150 |
 | f32-out GEMV audit | **CLAIM WRONG**: 35B runs 41 `CastF32`/step (3.1%) | Fold into the 35B lever |
 | Invocation-parity | CI guard + checklist landing | build-verify `kGemvHeuristicAlgos` |
 | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B |

.agents/roadmap_v1.md

@@ -44,6 +44,7 @@ issue is not yet placed. Keyed record: update in place, never append.
 | [#230](https://github.com/mudler/vllm.cpp/issues/230) | — | `test_agent_record.py`: 7 issue-intake tests never run, and error when they do | bug |
 | [#222](https://github.com/mudler/vllm.cpp/issues/222) | `SERVE-GATE-ONLINE` | The online-serving gate aborts on current main: harness looks for `examples/server`, the build produces `examples/vllm-server` — paths repaired by `2b262622`/`8fce04d3`; residue is the drift guard, the `27n` key and a model gate that recorded a skip as passed | bug |
 | [#224](https://github.com/mudler/vllm.cpp/issues/224) | `ENG-DOCS-SITE` | Publish `docs/` as a GitHub Pages site without owning a second copy | feature |
+| [#213](https://github.com/mudler/vllm.cpp/issues/213) | `PERF-27B-LMHEAD-FP4` | Qwen3.6 NVFP4 baselines (27B and 35B-A3B) must reach vLLM speed parity | perf |
 | [#203](https://github.com/mudler/vllm.cpp/issues/203) | `BACKEND-VULKAN` | Vulkan on unified memory holds TWO copies of the weights: 27B peaks at 100.8 GiB RSS and OOM-reboots a Spark | bug |
 | [#201](https://github.com/mudler/vllm.cpp/issues/201) | `BACKEND-ROCM` | `hipblasGemmEx` overload mismatch in `rocm_matmul_hipblaslt.hip` | bug |
 | [#199](https://github.com/mudler/vllm.cpp/issues/199) | `BACKEND-METAL-MLX` | macOS MLX build fails on `-Werror` in MLX headers | bug |

docs/BENCHMARKS.md

@@ -94,6 +94,20 @@ the same metric at higher concurrency (c8 p99 ITL 0.86x, but 1.055x at c16 and
 | Roof | 20.42 GiB over about 273 GB/s: vLLM's 79 ms/token is about 95% of the bandwidth limit, ours 96 ms/token about 78% | | | |
 | Peak host RSS | 21.0 GiB, down from 24.2 GiB, because the FP8 tower is no longer expanded to BF16 | | | |
 
+#### NVFP4 `lm_head` kept packed (`PERF-27B-LMHEAD-FP4`, #213)
+
+| Axis | Packed (`VT_LMHEAD_FP4=1`) | Dequant (`=0`) | Result |
+|---|---:|---:|---|
+| Peak host RSS | 19.36 GiB | 21.06 GiB | **-1.70 GiB**, but measured BEFORE `ENG-LOAD-DIRECT-UPLOAD` (#150) made `LoadCtNvfp4Raw` borrow mmap'd bytes, which moves the RSS accounting; re-measurement OWED |
+| Peak host RSS, non-CUDA | | | **Arithmetic, not measured.** A backend with no fp4 GEMM keeps packed + one bf16 operand: Vulkan **-1.70 GiB** (it used to stage a host bf16 head *and* a device copy), plain CPU **+0.67 GiB**. See `docs/USAGE.md` |
+| Greedy continuation | identical to the dequant leg, byte for byte | | SOLID |
+| `test_qwen27_paged_engine` | 235/235 | 235/235 | unchanged |
+| tok/s, leg A / leg B | 11.197 / 11.193 | 9.418 / 10.163 | **INDICATIVE ONLY** |
+| Reading, throughput | packed faster in all four legs; packed legs agree to 0.04%, dequant legs disagree by 7.9% | | DIRECTION established, MAGNITUDE not |
+| Owed | binding grid: 3 reps per leg, order-alternated, c1/c2/c4/c8, medians of per-rep medians, before any ratio is quoted | | PENDING |
+| Method | same model and revision as the table above, same-binary A/B on GB10 | | |
+| Provenance | the row's fresh reviewer, gate checkpoint `nvidia/Qwen3.6-27B-NVFP4`@`0893e1606ff3d5f97a441f405d5fc541a6bdf404` | | |
+
 ### Qwen3.6-35B-A3B by concurrency
 
 | Concurrency | 1 | 2 | 4 | 8 | 16 | 32 |

docs/ENVIRONMENT.md

@@ -90,6 +90,7 @@ portable/reference path. In normal operation leave them unset.
 | `VT_CONV_REG` | on (CUDA GDN) | The non-register-tiled short causal convolution |
 | `VT_CONV_EXACT_CHUNKS` | on (CUDA GDN prefill) | Use `=0` for the legacy sequence-serial causal-conv mapping; default mirrors vLLM's exact `(sequence, 8-token chunk)` descriptor and is byte-identical |
 | `VT_MODELOPT_W4A4` | `0` (Qwen3.6 dense ModelOpt NVFP4) | ModelOpt NVFP4 checkpoints ship a per-tensor `input_scale` next to every projection. Consuming it sets `Nvfp4Weight::alpha`, which flips `IsTrueW4A4()` and routes the weight to the fp4-ACTIVATION GEMM; on `nvidia/Qwen3.6-27B-NVFP4` that produced incoherent text, so the default leaves `alpha` at 0 and takes the W4A16 weight-only dispatcher (verified coherent). Set `1` to consume `input_scale` and take the W4A4 path |
+| `VT_LMHEAD_FP4` | **on** (Qwen3.6 dense NVFP4 `lm_head`) | Keeps a ModelOpt/compressed-tensors NVFP4 output head PACKED (`Qwen3_5DenseWeights::lm_head_fp4`) so the logits GEMM reads `K*N/2 + K*N/16` bytes per step instead of the `2*K*N` of a dequantized bf16 operand (~0.715 GB vs ~2.543 GB on `nvidia/Qwen3.6-27B-NVFP4`), and the operand keeps its on-disk `[N,K]` orientation instead of forcing the row-major NN GEMM that has no `nvjet_sm121` kernel. Mirrors vLLM, which resolves a quantized `lm_head` through `ModelOptNvFp4W4A16LinearMethod` (`modelopt.py:2491-2496,2508-2536`) and never materializes bf16 (`logits_processor.py:98-133`). Measured peak host RSS 21.06 -> 19.36 GiB (**-1.70 GiB**: the bf16 head is 2,543,206,400 B = 2.368 GiB, the packed head 715,264,000 B = 0.666 GiB) — measured before #150 changed the RSS accounting, so the figure is owed a re-measurement. On a backend with NO fp4 GEMM the head additionally keeps ONE dequantized `2*K*N` bf16 operand, built at prepare time; no other NVFP4 projection keeps one, so a quantized tower is never expanded (issue #203). The head is W4A16 under BOTH spellings: the on-disk activation divisor (`input_scale` / `input_global_scale`) is dropped for the head unless `VT_MODELOPT_W4A4=1`, because vLLM's `ModelOptNvFp4W4A16LinearMethod` deletes it (`modelopt.py:1365`). `=0` is the same-binary rollback to dequantize-at-load. BF16, FP8, GGUF and tied heads are unaffected either way (row `PERF-27B-LMHEAD-FP4`, issue #213) |
 | `VT_FA2_PREFILL` | on (CUDA) | The portable prefill attention instead of the vendored FA2 |
 | `VT_FA2_DECODE` | on (CUDA) | The portable decode attention instead of the vendored FA2 |
 | `VT_FA2_DECODE_4B` | on (CUDA, Qwen3.5-4B) | The portable paged decode attention instead of the ratio-4 vendored FA2 path; the 27B and 35B selectors are unchanged |

docs/FEATURES.md

@@ -65,6 +65,7 @@ are our reading of their documented behavior, not measurements.
 | Format | vllm.cpp | vLLM | SGLang | llama.cpp |
 |---|---|---|---|---|
 | NVFP4 (W4A4 and W4A16 Marlin) | ✅ | ✅ | ✅ | ☐ |
+| NVFP4 `lm_head` kept packed (no dequant at load) | ✅ `VT_LMHEAD_FP4` default-ON, #213; CUDA-gated on `nvidia/Qwen3.6-27B-NVFP4`@`0893e160` (continuations byte-identical packed vs dequant, `test_qwen27_paged_engine` 235/235; RSS -1.70 GiB owed a re-measure) | ✅ | ☐ | ☐ |
 | GGUF k-quants and i-quants | ✅ (CPU grouped keep-quant MoE took a bf16-activation regression in `b4f5610a`; found by bisect and fixed 2026-08-06) | ☐ | ☐ | ✅ |
 | AWQ | ◐ CPU dequant | ✅ | ✅ | ☐ |
 | GPTQ | ◐ CPU dequant | ✅ | ✅ | ☐ |

docs/STATUS.md

@@ -75,7 +75,7 @@ token-for-token correctness against the pinned oracle.
 
 | Capability | State | Notes |
 |---|---|---|
-| Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy |
+| Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` **0.85x BEHIND**; its NVFP4 `lm_head` PACKED (#213; RSS remeasure due) |
 | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF |
 | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed |
 | Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact GDN chunks default ON and byte-identical to rollback. Local A/B: total/output +2.152%, TTFT -2.945%, TPOT/ITL -1.920%; sealed-vLLM comparison 1.021x throughput, 1.086x TTFT, 1.025x TPOT, +233 MiB VRAM ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) |

docs/USAGE.md

@@ -285,6 +285,34 @@ independently selectable with `VT_CPU_Q8_DOT`, `VT_CPU_QUANT_MMLA`, and
 while an unavailable forced tier fails closed. The exact accepted values are
 listed in [ENVIRONMENT.md](ENVIRONMENT.md).
 
+### The NVFP4 output head
+
+On a Qwen3.6 dense checkpoint whose `lm_head` is stored NVFP4 (ModelOpt
+`weight`/`weight_scale`/`weight_scale_2`, or compressed-tensors
+`weight_packed`/`weight_global_scale`) the head is kept **packed** and the logits
+GEMM runs on it directly, as vLLM does. Nothing is dequantized at load, so the
+head costs `K*N/2 + K*N/16` bytes instead of `2*K*N`, about 0.715 GB instead of
+2.543 GB on `nvidia/Qwen3.6-27B-NVFP4` (measured peak host RSS 21.06 to 19.36
+GiB, a 1.70 GiB saving on CUDA; the figure is owed a re-measurement after
+`ENG-LOAD-DIRECT-UPLOAD` changed the RSS accounting).
+
+That accounting is CUDA's. A backend with no fp4 GEMM (CPU, Vulkan, Metal, HIP,
+Tenstorrent) has to multiply against a dequantized bf16 copy, so on those the
+head costs the packed bytes **plus** one `2*K*N` operand, built once when the
+model is prepared rather than per call — 0.666 + 2.368 = 3.034 GiB on the same
+checkpoint. The sign of the change therefore depends on the backend: on Vulkan,
+which used to stage a host bf16 head *and* a device copy of it, the head goes
+4.736 to 3.034 GiB, the same **-1.70 GiB**; on plain CPU it goes 2.368 to 3.034,
+a **+0.67 GiB** regression, paid once instead of rebuilding 2.368 GiB on every
+decode step as that backend did before. Only the head is kept that way; every
+other NVFP4 projection dequantizes per call, so a quantized tower is never
+expanded in memory. The head runs W4A16 under both namings: the on-disk
+activation divisor next to it (`input_scale`, or `input_global_scale` in the
+compressed-tensors spelling) is NOT consumed unless `VT_MODELOPT_W4A4=1`,
+matching vLLM, which deletes it on this path. Set `VT_LMHEAD_FP4=0` for a
+same-binary A/B that restores the old dequantize-at-load owner. BF16, FP8, GGUF
+and `tie_word_embeddings` heads are unaffected by either setting.
+
 ### Validating a staged release archive
 
 Release verification reads only a freshly extracted archive, never files from

include/vllm/model_executor/models/qwen3_5_dense.h

@@ -97,14 +97,24 @@ struct Qwen3_5DenseLayerWeights {
   DenseMlpWeights mlp;                   // every layer has a dense MLP
 };
 
-// Whole dense-model text weights. `lm_head` is always materialized bf16 here, but
-// the CHECKPOINT may store it BF16, FP8 (per-channel scale) or ModelOpt NVFP4 —
-// the 27B NVFP4 publishers disagree, and revisions of one repo disagree with each
-// other (issue #164). LoadLmHeadAnyDtype dequantizes all three to this operand.
+// Whole dense-model text weights. The CHECKPOINT may store the head BF16, FP8
+// (per-channel scale) or ModelOpt NVFP4 — the 27B NVFP4 publishers disagree, and
+// revisions of one repo disagree with each other (issue #164). BF16/FP8 are
+// materialized into `lm_head`, NVFP4 stays PACKED in `lm_head_fp4`
+// (PERF-27B-LMHEAD-FP4, issue #213); exactly one is populated.
 struct Qwen3_5DenseWeights {
   OwnedTensor embed_tokens;  // bf16 [vocab, H]  (NOT transposed; embed lookup)
   OwnedTensor final_norm;    // bf16 [H]
   OwnedTensor lm_head;       // bf16 [H, vocab]  (dequantized -> Matmul-B layout)
+  // NVFP4-resident output head [N=vocab, K=H], kept in the on-disk orientation the
+  // fp4 GEMMs read. Mirrors Qwen3_5MoeWeights::lm_head_fp4 and vLLM's own decision
+  // to leave the head quantized: get_quant_method accepts ParallelLMHead
+  // (modelopt.py:2508-2536) over the bare `lm_head` key (modelopt.py:2491-2496),
+  // so ModelOptNvFp4W4A16LinearMethod — pinning MarlinNvFp4LinearKernel
+  // (modelopt.py:1249,1283-1284) — resolves it and logits_processor._apply_head
+  // calls quant_method.apply every step (logits_processor.py:98-133). Empty on
+  // every BF16/FP8/GGUF/tied checkpoint.
+  Nvfp4Weight lm_head_fp4;
   // Mirrors tie_word_embeddings: logits reuse embed_tokens as raw [V,H]
   // torch-Linear storage, so no second host/device owner is created.
   bool tied_lm_head = false;
@@ -135,6 +145,25 @@ OwnedTensor LoadLmHeadAnyDtype(const TensorResolver& get,
                                const std::function<bool(const std::string&)>& has,
                                const std::string& name);
 
+// PERF-27B-LMHEAD-FP4 (issue #213). Load the dense output head into EXACTLY ONE
+// of the two owners: a ModelOpt/compressed-tensors NVFP4 head stays PACKED in
+// `fp4_out`, every other storage form is materialized bf16 [in, out] into
+// `bf16_out` by LoadLmHeadAnyDtype. `proj` omits the trailing ".weight".
+void LoadDenseLmHead(const TensorResolver& get,
+                     const std::function<bool(const std::string&)>& has,
+                     const std::string& proj, OwnedTensor& bf16_out,
+                     Nvfp4Weight& fp4_out);
+
+// True when the checkpoint ships an EXPLICIT head under either naming
+// (`<proj>.weight`, or `<proj>.weight_packed` for compressed-tensors NVFP4);
+// false means `tie_word_embeddings`.
+bool DenseCheckpointHasLmHead(const std::function<bool(const std::string&)>& has,
+                              const std::string& proj);
+
+// VT_LMHEAD_FP4 (default ON): the in-binary rollback for the packed head. `0`
+// restores the dequantize-at-load owner, so the A/B is same-binary.
+bool DenseLmHeadFp4Enabled();
+
 Fp8Weight LoadFp8RawShared(const TensorResolver& get, const std::string& proj);
 
 OwnedTensor MaterializeCtNvfp4Bf16Transposed(const TensorResolver& get,
@@ -161,6 +190,17 @@ Qwen3_5DenseLayerWeights LoadQwen3_5DenseLayer(const TensorResolver& get,
                                                const std::string& layer_type,
                                                int64_t layer_idx);
 
+// The same load with an EXPLICIT presence probe — what `LoadQwen3_5Dense` calls
+// per layer. The resolver-only overload above answers `has` with a constant
+// `true`, which forces every routed projection down the compressed-tensors
+// spelling; a checkpoint that mixes namings (the ModelOpt `weight_scale_2` form,
+// or an FP8/BF16 projection next to an NVFP4 one) needs the real probe. Exposed
+// so the loader gate can drive a whole synthetic layer through the SAME routing
+// production takes.
+Qwen3_5DenseLayerWeights LoadQwen3_5DenseLayer(
+    const TensorResolver& get, const std::function<bool(const std::string&)>& has,
+    const std::string& layer_type, int64_t layer_idx);
+
 // Full dense-model load across the given shards. Uses config.num_hidden_layers
 // and config.layer_types. Text path only — the vision tower (model.visual.*)
 // and image/video merger are DEFERRED (notes §0.1). The checkpoint's MTP
@@ -188,6 +228,13 @@ class Qwen3_5DenseModel {
   static void PrepareBf16Resident(const Qwen3_5DenseWeights& weights,
                                   vt::Queue& queue);
 
+  // PERF-27B-LMHEAD-FP4 (issue #213). Build the resident form of the packed
+  // `lm_head_fp4` THIS backend's logits GEMM consumes: the Marlin W4A16 repack
+  // on CUDA (PRE-CAPTURE), else the dequantized bf16 [K,N] operand. Inert when
+  // the head is not packed. Called from the registry `prepare` hook.
+  static void PrepareLmHeadResident(const Qwen3_5DenseWeights& weights,
+                                    vt::Queue& queue);
+
   // Batched PAGED dense forward — the 27B analogue of Qwen3_5Model::Forward.
   // Same signature/structure (paged KV cache for the full-attn layers, batched
   // GDN recurrent state for the GDN layers, the f32 residual thread), reusing the

include/vllm/model_executor/models/qwen3_5_weights.h

@@ -238,6 +238,16 @@ struct Nvfp4Weight {
   // path. Uploaded once from the persistent `alpha` member; the diagnostic host
   // scalar path leaves this null.
   mutable std::shared_ptr<void> d_alpha;
+  // OPT-IN lifetime residency for the DEQUANTIZED bf16 [K=in, N=out] Matmul-B
+  // operand the backends with NO fp4 GEMM multiply against (CPU / Vulkan / Metal /
+  // HIP / Tenstorrent; CUDA never dequantizes). Default OFF, and it must stay a
+  // per-WEIGHT opt-in: the operand is a bf16 expansion of ~4x the packed bytes, so
+  // holding one per tower projection is the double-residency that OOM-reboots a
+  // Spark on Vulkan (#203). The dense loader opts in the OUTPUT HEAD alone — one
+  // weight, re-read whole every step (~2.54 GB a step rebuilt per call at the
+  // 27B's 248320x5120); everything else keeps a per-call copy.
+  bool keep_dequant_b = false;
+  mutable std::shared_ptr<void> d_dequant_b;
 
   // Resident Marlin constants (issue #237; see ResidentSlot). `resident_marlin`
   // is this weight's own repack; `resident_marlin_pair` is the fused gate+up