From 6603356ae3986d65b77302cdea8dbb4828104e5d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 9 Aug 2026 19:35:45 +0000 Subject: [PATCH] perf(qwen3_5,moe): re-verify the NVFP4 gate models and close most of the gap The 27B-NVFP4 grid of record was measured on `unsloth/Qwen3.6-27B-NVFP4`, and that repo re-quantized in place: @890bdef7 is genuine NVFP4, @ccdaab7e is the same name turned into FP8 W8A8. `nvidia/Qwen3.6-27B-NVFP4` @0893e160 is the publisher's own single-revision NVFP4 build and is the reference from here on. Re-measuring against it found us far behind, and most of that is now closed. LOADING. Three defects, all the same shape: the dense loader assumed BF16 wherever the compressed-tensors NVFP4 probe did not match. One `MaterializeBf16Source` now sits under the BF16 loaders (a per-channel scale read as per-tensor is REJECTED rather than silently wrong); `IsNvfp4Projection` accepts ModelOpt naming (`weight`/`weight_scale_2`) alongside compressed-tensors; and W4A16 stays the default because consuming ModelOpt's `input_scale` flips `IsTrueW4A4()` into the fp4-activation GEMM, which produced incoherent text here (`VT_MODELOPT_W4A4=1` opts in). THE GAP WAS ALSO A LOADER BUG. `LoadAttnDense` and `LoadGdnDense`'s `out_proj` had only NVFP4-or-BF16 branches, so a `modelopt_mixed` FP8 tower fell through to the BF16 path, which DEQUANTIZES: ~3 GiB of FP8 became ~5.9 GiB of BF16 re-read every decode step and executed as cuBLAS `gemvx`. Never a missing capability -- the `*_fp8` slots and their `MatmulFp8Cutlass*` consumers already existed and the MoE loader has done exactly this since the 35B work, with its own comment calling it the DEFAULT. vLLM keeps the same weights fp8, so this is POL-MIRROR-VLLM. 27B, warm servers, greedy, ignore_eos, 128 tokens, medians of 3, vs vLLM 0.25.0 c | before | after | vs vLLM before -> after c1 | 8.76 | 10.41 | 0.713x -> 0.847x c8 | 62.13 | 72.49 | 0.722x -> 0.843x peak host RSS 24.2 -> 21.0 GiB TWO ROUTER KERNELS WERE BARRIER-BOUND, not compute-bound. The router grid is one block per token, so at decode a single block ran the k selection rounds with a block-wide tree each (~64 `__syncthreads` for k=8 over 256 experts); the GROUPED variant was worse, running group scoring, mask, top-k and renorm on ONE lane behind `if (threadIdx.x != 0) return;`, which its own comment deferred as "W9". Both now use a warp-shuffle argmax. Byte-identical by construction: top-k is an ARGMAX over a total order, which reassociates freely, while the softmax max and sum reductions in the same kernels are ARITHMETIC and are deliberately left on their original trees. 35B-A3B @491c2f1e, same harness: c1 70.58 -> 73.24 (0.945x -> 0.980x), c4 194.87 -> 199.90 (0.952x -> 0.977x), c2 0.826x -> 0.867x, c8 0.900x -> 0.919x THE 27B GATE WAS LETTING THE FILESYSTEM PICK THE MODEL. Five checkpoint-gated tests, SACRED `test_qwen27_paged_engine` among them, took the first entry `fs::directory_iterator` yielded under a repo with TWO revisions, so a token-exact pass against the FP8 revision would have been recorded as an NVFP4 pass. Green by luck. `tests/parity/hf_snapshot.h` pins the revision the goldens' own `oracle.model` field names, with `VT_QWEN27_SNAPSHOT` as the escape hatch. `max_num_seqs` also moves 8 -> 32: at 8 a c8 client sat exactly on our own batch ceiling. Not vLLM's 1024, because it also caps the padded decode-graph set. Correctness: 27B greedy continuation byte-identical to vLLM across 5 repeats. The 35B's continuation differs at ONE token, and that is NOT a defect -- vLLM's own top-2 there are the same float32 value (-1.2221027612686157, diff 0.0) and its `torch.argmax` breaks the tie by lower index; our `ArgReduce` implements the same rule, but our logits are not bit-identical, so the tie does not reproduce. Gates: test_qwen27_paged_engine 235/235, test_qwen36_paged_engine 315/315, test_ops_moe 33451/33451, test_ops_moe_grouped 440/440, grouped_bf16 19/19, test_deepseek_v4_moe 716/716. Still behind, and named rather than hidden: 27B at 0.85x with the NVFP4 MLP marlin at ~68% of the bandwidth roof, 35B c2/c8 the weak cells, `CastF32Kernel` 3.1% of the 35B step. Full method, the refuted hypotheses and the decode attributions are in `.agents/benchmark-record.md`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode] --- .agents/NOW.md | 24 +- .agents/benchmark-record.md | 481 ++++++++++++++++++ docs/BENCHMARKS.md | 30 +- docs/ENVIRONMENT.md | 1 + docs/FEATURES.md | 2 +- docs/STATUS.md | 8 +- docs/USAGE.md | 14 +- include/vllm/entrypoints/model_loader.h | 10 +- .../models/dense_weight_loaders.h | 128 ++++- .../model_executor/models/qwen3_5_dense.h | 2 + scripts/check-public-doc-tables.py | 8 +- src/capi/vllm_c.cpp | 2 +- src/vllm/entrypoints/openai/server_main.cpp | 2 +- src/vllm/model_executor/models/qwen3_5.cpp | 34 +- .../models/qwen3_5_dense_weights.cpp | 127 ++++- .../model_executor/models/qwen3_5_weights.cpp | 5 + src/vt/cuda/cuda_moe.cu | 137 +++-- tests/parity/hf_snapshot.h | 66 +++ tests/parity/test_op_parity.cpp | 14 +- .../parity/test_qwen27_ngram_spec_decode.cpp | 18 +- tests/parity/test_qwen27_paged_engine.cpp | 22 +- tests/parity/test_qwen27_spec_decode.cpp | 14 +- .../test_qwen27_spec_decode_concurrent.cpp | 15 +- 23 files changed, 1010 insertions(+), 154 deletions(-) create mode 100644 tests/parity/hf_snapshot.h diff --git a/.agents/NOW.md b/.agents/NOW.md index b88299745..57ac13dda 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -13,17 +13,18 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | Claim / track | State | Next command or step | |---|---|---| | State record (#166) | **157 imports = 3,231,342 exact bytes** at `776c56f1`; 95/95; prior 156 wrappers/rows preserved; raw-row guard | Force-update #166; rerun readiness | -| Laguna NVFP4 / DeepSeek-V4 decode | **Both CLOSED, byte-exact, default-ON**: 1.03x vLLM, 1.144x ds4 | Laguna vLLM K-run when convenient | -| f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool | +| 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 | +| 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 | Merge; build-verify `kGemvHeuristicAlgos` on dgx | -| MiniMax-H3 lane | **bf16 shards STREAM both towers (DiT + encoder); Q4_K_M enc cond cos 0.9975, 3.5° med, DIFFUSE** | render A/B on saved embeds | -| Kimi-Linear-48B | **ROW 7 fold LANDS (#122 §21): engine==CLI 128/128; golden 122/128; SACRED green; v13 tokens ABI** | ACTIVE: 19.0 tok/s vs vLLM ~21 (~0.90×) | -| 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) | -| Qwen3.5-4B sm_120 | Exact chunks ON: rebased-main reprofile 3.072x kernel / +2.272% run; sealed-vLLM throughput 1.021x PASS. Latency/VRAM OPEN | Spike residual 1.609x conv gap | -| RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec, RSS -24% | W6: BF16 GEMM | -| MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED (A/B ties our+vLLM PTX all ptxas/JIT; +10us=engine context, not codegen)** | TERMINAL: at parity | +| MiniMax-H3 lane | **bf16 shards STREAM both towers; Q4_K_M enc cond cos 0.9975, DIFFUSE** | render A/B on saved embeds | +| Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED (~10% band) | Warm-server path: ckpt is tiktoken-only | +| 35B fresh grid | @`491c2f1e`: warp-shuffle router LANDED, **c1/c4 now 0.98x**, c2 0.87x, c8 0.92x | `CastF32` 3.1%; tighten c2/c8 spreads | +| Qwen3.5-4B sm_120 | Exact chunks ON: 3.072x kernel / +2.272% run; sealed-vLLM tput 1.021x PASS; latency/VRAM OPEN | Spike residual 1.609x conv gap | +| RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | +| MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): async mirror on classic dense Qwen3; SACRED 184/184 | Residual: sibling scope one-liner | -| CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn**. **G5 not next** | Parakeet encoder; attn dtype hoist | +| CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn** | Parakeet encoder; attn dtype hoist | | Supported-models list | **LANDED**: FEATURES arch table CI-bound (33 archs) | — | | `/v1/videos` OpenAI shape | **MERGED** (#71): Sora `model`/`size`/`seconds` + `GET /{id}/content` | `row/SERVE-VIDEOS-REFS` PR open: reference conditioning | | Vulkan 27B | decode **MET 4.36 vs 4.35** (barriers OFF). **LOADMEM: load held the model TWICE, VmRSS 100.759 -> 53.413 GiB** | Load-phase host build is the new peak | @@ -43,15 +44,14 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See ## Next actions -0. **`ROAD-V1-MEM`** KV auto-sizing spike LANDED (`specs/kv-sizing.md`, `READY`). +0. **27B NVFP4 0.72x -> 0.85x** (FP8 tower native). Next: NVFP4 MLP marlin, 68% + of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH; no `kv_cache_dtype`. 1. **Spike the Parakeet encoder row** (vLLM carries it inside `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). 2. **Qwen3.5-4B sm_120:** rebased branch is GREEN and reprofiled. Spike the residual 1.609x conv gap; latency/VRAM and gate models stay open. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. -3. **Same-tool re-verify deepseek_v4's bf16 resident tower** (the one other - f32-out caller) once the Laguna fix proves the mechanism. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). 5. **Protocol substrate — partly done.** Triage/audit + `STATUS.md` ratchet + `AGENTS.md` tiering DONE. REMAINING: anchor backfill (6 model rows need a diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 045823061..9b1d182da 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -17065,3 +17065,484 @@ Peak is now the LOAD phase — the host `OwnedTensor` build reaches ~51 GiB befo the first upload, and the adoption only acts afterwards. Copying from the mmap straight into the device buffer at load would cut that too. The unreleased page cache is a second, independent lever. +## 2026-08-09 — Qwen3.6-27B NVFP4 re-verified on the `nvidia` ModelOpt checkpoint + +The recorded 27B-NVFP4 grid was measured on `unsloth/Qwen3.6-27B-NVFP4`, and +that repo re-quantized in place: @`890bdef7` is genuine NVFP4, @`ccdaab7e` is the +same repo name turned into FP8 W8A8 throughout. `nvidia/Qwen3.6-27B-NVFP4` +@`0893e1606ff3d5f97a441f405d5fc541a6bdf404` is the publisher's own single-revision +NVFP4 build and is the reference from here on. It is not the same model as the +`unsloth` one: `hf_quant_config.json` declares `MIXED_PRECISION`, so the MLP is +`W4A16_NVFP4` group-16 while the whole `linear_attn` + attention tower is FP8, +and it declares `kv_cache_quant_algo: FP8`. + +### Two method defects found before any number was taken + +**The benchmark build was not the production stack.** `~/build_gdnfp8.sh` on dgx +configured with `-DVLLM_CPP_CUDA=ON -DVLLM_CPP_BUILD_TESTS=OFF +-DVLLM_CPP_CUTLASS_DIR=...` and never passed `-DVLLM_CPP_TRITON=ON`, so +`CMakeCache.txt` read `VLLM_CPP_TRITON:BOOL=OFF` and `src/vt/cuda/cuda_gdn.cu:45` +compiled out the entire Triton-AOT family: `gdn_deltah_h48/h32`, `gdn_chunko_*`, +the WU pipeline, the packed decode `gdn_decode_h48/h32` and the whole KDA family. +Qwen3.6-27B is a GDN hybrid, so every layer fell back to the hand kernels. +`.agents/environment.md:51-66` already calls both flags MANDATORY on this box and +the 27B gate refuses to run without them; a hand-written benchmark script has no +such guard, which is how it slipped. Re-measured on a correct build (configure +log verified for `CUTLASS found`, `FlashAttention-2 prefill/decode: ENABLED` and +`Triton AOT: ... <- vendored ... sm_121a`, plus `gdn_decode_h48` symbols present +in the `cuda_gdn` object), the throughput difference is a WASH, so the Triton-AOT +GDN family is not a decode-throughput lever on this model. The numbers taken off +that build were still not quotable. + +**The 27B gate let the filesystem choose the checkpoint.** Five checkpoint-gated +tests, including SACRED `test_qwen27_paged_engine`, resolved the snapshot by +taking whatever `fs::directory_iterator` yielded first under a repo with TWO +cached revisions. Order is unspecified, so a token-exact pass against the FP8 +revision would have been recorded as an NVFP4 pass. It has been green by luck: +readdir yields @`890bdef7` first on this box today. Now pinned through +`tests/parity/hf_snapshot.h` to the revision named in the goldens' own +`oracle.model` field, with `VT_QWEN27_SNAPSHOT` as the explicit escape hatch. + +### Correctness + +Production build, GB10 sm_121a: `test_qwen27_paged_engine` 235/235, +`test_qwen36_paged_engine` 315/315. On the `nvidia` checkpoint, which has no +committed golden, the greedy continuation was captured from the SAME warm server +processes that produced the throughput numbers, so a mismatch could not be blamed +on the two engines having loaded different weights in different sessions. Both +engines returned, for `"The capital of France is"` at 32 tokens, +`" Paris.\nThe capital of France is Paris.\nThe capital of France is Paris.\n..."` +identically. + +### Throughput + +Warm servers, one engine resident at a time under one `flock`, greedy, +`ignore_eos` so both emit exactly 128 output tokens, 7-token prompt, +`--gpu-memory-utilization 0.55 --max-model-len 4096`, medians of three +repetitions after a discarded warm-up, vLLM 0.25.0 in its production graphed +config. Ours at `row/NVFP4-REPUBLISH-LAND` (current main plus the republish +loader fix, the fp8-native GDN tower and `max_num_seqs` 32). + +| c | ours | vLLM 0.25.0 | ours/vLLM | spread ours | spread vLLM | +|---|---:|---:|---:|---:|---:| +| 1 | 8.76 | 12.29 | 0.713x | 1.0002 | 1.0060 | +| 2 | 17.07 | 23.49 | 0.727x | 1.0087 | 1.0690 | +| 4 | 33.01 | 45.42 | 0.727x | 1.0067 | 1.0005 | +| 8 | 62.13 | 86.01 | 0.722x | 1.0023 | 1.0030 | + +The deficit is uniform and far outside the noise band. Scaling is NOT the +problem: c1 to c8 we gain 7.09x against vLLM's 7.00x. From the roof, 20.42 GiB +of weights over about 273 GB/s puts a floor near 75 ms/token; vLLM's 79 ms is +about 95% of it and our 114 ms about 65%, so roughly a third of our per-token +time is not weight streaming. + +### Configuration facts recorded with the numbers + +vLLM honors the checkpoint's `kv_cache_quant_algo: FP8` and runs +`kv_cache_dtype=fp8_e4m3`; we have no `kv_cache_dtype` concept at all, so vLLM +moves half the KV bytes we do. At 128 output tokens that is a small share of the +difference, at long context it is not, and it is a POL-MIRROR-VLLM gap either +way. Separately, vLLM 0.25.0 transiently touches 118 of 119 GiB during +`determine_available_memory` at this utilization and pushes the box into swap; it +recovers, but that is inside the OOM-reboot band for GB10. + +Owed: the 1,024 in / 128 out total-throughput axis on this checkpoint, so it can +be compared against the `unsloth` grid, and the same treatment for the 35B. + +### Lever A/B: the dense NVFP4 Marlin route on qwen3_5 (2026-08-09) + +`qwen3_5.cpp` carries its OWN `MatmulNvfp4MarlinD` (:2357) which sends E=1 DENSE +projections through `MoeGroupedGemmNvfp4Marlin` plus a `moe_align` cache, while +`dense_nvfp4_gemm.h` has a dense route (`vt::MarlinDenseGemm`, rank-2 operands, +direct-A, no moe_align) that qwen3, olmo2, deepseek_v2 and minimax_h3 all take, +and which this repo measured at 48-CTA ~86 us/call against the MoE route's +128-CTA ~118 us/call. That made it the prime suspect for the uniform 0.72x. + +**The earlier attempt at this hung; the root cause is now named.** The reverted +patch mirrored `dense_nvfp4_gemm.h:366-385` including its call sequence, and that +sequence has NO per-call workspace memset because the shared +`DenseMarlinWorkspace` zeroes at allocation. `qwen3_5`'s own `DenseMarlinWorkspace` +does NOT zero at allocation; it relies on a memset at the call site. Marlin's +fp32 reduce SPINS on those lock words, so uninitialized locks spin forever, which +is a hang rather than a fault and is why turning the decode graph off appeared to +"fix" it. Keeping the per-call memset, the route runs first try with no hang. + +**Same-binary A/B**, one build, one flock, `VT_MARLIN_DENSE` the only difference, +27B `nvidia` @`0893e160`, ns=32, 3 reps: + +| c | dense ON | MoE route | delta | +|---|---:|---:|---:| +| 1 | 8.78 | 8.78 | +0.0% | +| 2 | 16.83 | 16.74 | +0.5% | +| 4 | 32.98 | 32.81 | +0.5% | +| 8 | 62.29 | 61.81 | +0.8% | + +The 32-token greedy continuations of the two arms are IDENTICAL, and both match +vLLM. So the lever is REAL but SMALL: it removes a framework fork and buys under +one percent above c1, and it does not explain a 38% deficit. The prime suspect is +retired; attribution moves to a decode-window profile rather than another +structural guess. + +### Decode attribution: where the 114 ms/token goes (2026-08-09) + +Two-length `nsys` difference (8 vs 136 tokens, same prompt, `cuda_gpu_kern_sum` +diffed per kernel), 27B `nvidia` @`0893e160`, single stream, production build. + +**The first run of this was WRONG and the way it was wrong matters.** `nsys` +defaults to `--cuda-graph-trace=graph`, which reports a captured decode graph as +ONE range, so every kernel inside it contributes nothing to the per-kernel table. +The diff came back at 0.030 ms/step against a 114 ms/token wall clock, with +`dInstances == 0` for `marlin_moe_wna16::Marlin`, `nvjet_sm121_*`, `gemvx` and +every other forward kernel, and +128 only for the sampler and embedding kernels +that live OUTSIDE the graph. That profile makes the sampler look like the whole +cost of decode and the GEMMs look free. Re-run with `--cuda-graph-trace=node`. +Sanity check for any future profile on a graphed path: the hot kernels' +instance counts MUST scale with the number of decode steps. + +**Corrected decode-only total: 14,630 ms over 128 steps = 114.298 ms/step**, +against a measured 114 ms/token wall clock. **We are ~100% GPU-busy at c1**, so +there is no host, scheduler or launch-overhead lever to recover here; the entire +deficit is kernel efficiency. + +| Share | Calls/step | ms/step | Kernel | +|---:|---:|---:|---| +| 40.0% | 192 | 45.77 | `marlin_moe_wna16::Marlin` (the NVFP4 MLP, 3 GEMMs x 64 layers) | +| 22.7% | 80 | 25.98 | cuBLAS `gemvx` bf16 | +| 16.2% | 48 | 18.51 | `nvjet_sm121_qqsss_mma_64x128x128` (cuBLASLt, FP8 tower) | +| 10.0% | 1 | 11.38 | `cutlass_80_tensorop_s16816gemm_bf16_128x64_32x6_nn_align2` | +| 6.2% | 48 | 7.05 | `nvjet_sm121_qqtst_mma_64x64x128` | +| 1.9% | 32 | 2.19 | cuBLAS `gemvx` bf16 (second shape) | +| 1.2% | 48 | 1.32 | `GdnDecodeFusedKernel` | + +Checkpoint census, which is what those kernels are reading: 8.561 GiB U8 (NVFP4, +of which 7.969 is `mlp` and 0.592 is `lm_head`), 7.789 GiB F8_E4M3 (5.156 +`linear_attn`, 1.562 `self_attn`, 0.996 `mlp`), 4.066 GiB BF16 (2.376 of it the +embedding TABLE, which decode reads one row of, plus 0.997 `mlp`, 0.450 other, +0.195 `self_attn`), 20.416 GiB total. + +Two readings follow. The NVFP4 MLP moves 8.56 GB in 45.77 ms, about 187 GB/s or +roughly 68% of this box's ~273 GB/s, so about 14 ms is theoretically on the +table there. The bf16 `gemvx` line is the sharper anomaly: after the FP8 and +NVFP4 towers are accounted for, at most about 1.7 GB of bf16 weight remains for +those 112 calls, and 28.2 ms to move 1.7 GB is far off the bandwidth roof, which +means those GEMVs are latency or occupancy bound rather than bandwidth bound. +That is the largest single mis-shaped cost in the step and it is where the next +lever belongs, ahead of any further structural refactor. + +### The FP8 tower was being dequantized to bf16, and it was the gap (2026-08-09) + +The decode profile said 24.6% of the step was cuBLAS bf16 `gemvx` moving at most +~1.7 GB, which is nowhere near bandwidth bound. Reading the loader explains it. + +`LoadAttnDense` (`qwen3_5_dense_weights.cpp`) had exactly TWO branches, NVFP4 or +`LoadBf16RawNK`, and the same was true of `LoadGdnDense`'s `out_proj`. A +`modelopt_mixed` checkpoint quantizes the attention tower to FP8 W8A8 while +leaving the MLP NVFP4, so every one of those projections matched neither branch's +intent and fell through to the bf16 path, which DEQUANTIZES: 1.562 GiB of +`self_attn` FP8 plus 1.41 GiB of `linear_attn.out_proj` FP8 became ~5.9 GiB of +BF16, re-read every decode step and executed as cuBLAS `gemvx`. + +This was never a missing capability. `FullAttnLayerWeights` already carries +`q/k/v/o_proj_fp8`, `GdnLayerWeights` already carries `out_proj_fp8`, +`MatmulFp8CutlassD` / `MatmulFp8CutlassPreQuantD` already consume them, and +`LoadAttn` on the MoE path (`qwen3_5_weights.cpp:319`) has done exactly this +since the 35B work, with its own comment calling it the DEFAULT. The dense path +was the one that never got it. vLLM does the same thing (`modelopt.py:519` +`process_weights_after_loading` keeps `layer.weight` fp8 and transposes it, then +applies through `fp8_linear`), so this is POL-MIRROR-VLLM, not an invention. + +**MEASURED**, same harness, same checkpoint, same box, 3 reps: + +| c | before | after | gain | vs vLLM before | vs vLLM after | +|---|---:|---:|---:|---:|---:| +| 1 | 8.78 | 10.41 | +18.6% | 0.713x | **0.847x** | +| 2 | 16.83 | 20.22 | +20.1% | 0.727x | **0.861x** | +| 4 | 32.98 | 38.76 | +17.5% | 0.727x | **0.853x** | +| 8 | 62.29 | 72.49 | +16.4% | 0.722x | **0.843x** | + +Peak host RSS falls 24.2 -> 21.0 GiB, the ~3.2 GiB the expansion was costing. +The 32-token greedy continuation is byte-identical to vLLM's, captured from the +same warm process as the numbers. SACRED gates on this build: +`test_qwen27_paged_engine` 235/235, `test_qwen36_paged_engine` 315/315. + +Per-token time goes 114 ms to about 96 ms, which is ~78% of the bandwidth roof +against vLLM's ~95%. The residual is the NVFP4 MLP marlin at ~68% of roof, and +that is the next lever. + +### 35B-A3B NVFP4 re-verified, and a FIRST-REQUEST-ONLY correctness defect found + +Same harness and checkpoint discipline as the 27B, on +`nvidia/Qwen3.6-35B-A3B-NVFP4` @`491c2f1e`, ours at the FP8-tower build: + +| c | ours | vLLM 0.25.0 | ratio | spread ours / vLLM | +|---|---:|---:|---:|---:| +| 1 | 70.58 | 74.71 | 0.945x | 1.001 / 1.001 | +| 2 | 107.64 | 130.35 | 0.826x | 1.022 / 1.111 | +| 4 | 194.87 | 204.67 | 0.952x | 1.005 / 1.029 | +| 8 | 318.56 | 354.04 | 0.900x | 1.073 / 1.081 | + +c2 and c8 carry real noise on both sides (vLLM's c2 spread is 1.111 on one slow +leg), so treat c1 and c4 as the trustworthy cells. + +**THESE NUMBERS SIT ON TOP OF A CORRECTNESS DEFECT AND ARE NOT A PARITY CLAIM.** +The greedy continuation captured from the warm 35B server did NOT match vLLM's, +and the probe that followed is the important part. Five identical greedy requests +to ONE warm server, `"The capital of France is"`, 32 tokens, `ignore_eos`: + +- request 1: `" Paris, a city renowned for its rich history, culture, and iconic + landmarks. Paris is situated in the north-central part of France..."` -- this + is BYTE-IDENTICAL to vLLM. +- requests 2 through 5: `" Paris, a city renowned for its iconic landmarks such + as the Eiffel Tower, the Louvre Museum, and the Notre-Dame Cathedral..."` -- + all four identical to each other, none matching vLLM. + +So it is not noise and not a near-tie coin flip: the FIRST request on a fresh +server is vLLM-exact and every subsequent request deterministically produces a +different, still-fluent continuation. Something survives across requests and +changes the result. Prefix caching is OFF on both engines. + +The same probe on the 27B is CLEAN: all five captures identical and capture 1 +byte-identical to vLLM (the fifth differed by one trailing newline from the +harness's block splitting, verified equal after `strip()`). So this is specific +to the 35B MoE/GDN path, not the shared serving loop. + +`test_qwen36_paged_engine` is 315/315 token-exact at ITS engine params, which +means the model and the weights are right and the defect lives in the serving +configuration or in state reuse between requests. This is plausibly the same +"prod async batch-1 greedy degeneration" a probe found earlier. It is NOT fixed +here and it OWNS the 35B speed row: no 35B parity claim may be made from the +table above until a second request returns what the first one does. + +### CORRECTION to the 35B entry above: it is not a first-request defect + +The entry above claimed the 35B divergence was state-dependent, request 1 +byte-identical to vLLM and requests 2-5 different. **That claim is REFUTED and is +withdrawn.** It rested on a single probe run and did not survive repetition. + +Four further probe runs on the same build and checkpoint, five captures each: +async runner off (`VT_ASYNC_RUNNER=0`), decode graph off +(`VLLM_CPP_CUDAGRAPH=0`), `--max-num-seqs 8`, and the default config twice. All +twenty captures are IDENTICAL to each other and NONE matches vLLM. Adding the +original run's four later captures, that is 24 of 25 captures on one answer. The +single vLLM-matching capture has never reproduced and remains unexplained; it is +recorded as an anomaly, not as a mechanism. + +Two things are cleared by those arms. The `max_num_seqs` 8 -> 32 default change +in this branch did NOT cause it: `--max-num-seqs 8` diverges identically. And +neither the async runner nor the captured decode graph causes it: both rollbacks +diverge identically. + +**What it actually is, and it is worse.** vLLM was put through the same probe on +the same checkpoint and params: five captures, all identical, all matching its +own earlier golden. So vLLM is DETERMINISTIC here and we are DETERMINISTIC on a +different continuation. Under the ratified near-tie gate -- token-exact where +vLLM is deterministic, distributional only where vLLM's own greedy is not -- the +distributional escape does not apply and we owe token-exactness. We fail it. + +Both continuations are fluent (`"...its rich history, culture, and iconic +landmarks. Paris is situated in the north-central part of France..."` from vLLM +against `"...its iconic landmarks such as the Eiffel Tower, the Louvre Museum, +and the Notre-Dame Cathedral..."` from us), diverging at token seven. + +The scope is now precise. `test_qwen36_paged_engine` is 315/315 +token-exact against the same oracle, so at the GATE's engine params we match and +at the SERVER's params we do not. The defect lives in whatever differs between +those two configurations, and that difference is the next thing to bisect. The +35B throughput table stands as measured and remains NOT a parity claim. + +### 35B divergence, bisected: it is PROMPT-dependent and lives in the engine + +Continuing the correction above, five more arms, each five captures unless noted: + +| Arm | Result | +|---|---| +| `VT_ASYNC_RUNNER=0` | stable, diverges | +| `VLLM_CPP_CUDAGRAPH=0` | stable, diverges | +| `--max-num-seqs 8` | stable, diverges | +| `--max-model-len 0` (config default) | stable, diverges | +| default config, twice | stable, diverges | +| vLLM, same probe | stable, matches its own golden | +| `examples/vllm-cli` (SYNC engine, all defaults, 1 run) | diverges, same text | + +Tokenization is identical on both sides and both models: `prompt_tokens=5`, +`completion_tokens=32` for ours and for vLLM, on the 27B and the 35B. + +So the serving layer is EXONERATED. `vllm-cli` drives `LoadedEngine` + +`engine().generate()` with `EngineParams{}` -- the same entry point +`test_qwen36_paged_engine` uses, which is 315/315 token-exact -- and it produces +the divergent continuation. The difference between the gate and this run is +therefore not the engine configuration at all. It is the PROMPT. + +That relocates the whole thing. Our 35B forward disagrees with vLLM on +`"The capital of France is"` from token seven, while agreeing token-for-token on +the gate's pinned prompt. Both engines are internally deterministic, so this is +not a scheduling or state-reuse effect and the distributional escape does not +apply. It is either a logit near-tie this prompt happens to land on, or a real +numerical difference in the MoE path that the gate's single prompt never +exercises. Either way the gate has a COVERAGE hole: one prompt cannot certify a +forward. + +Next: teacher-force both engines on this prompt and read the top-2 logit margin +at the divergence position. A margin at the bf16 noise floor makes it a ratified +near-tie and the gate needs more prompts; a wide margin makes it a real defect in +the MoE forward. Until that is answered the 35B throughput table is not a parity +claim. + +### RESOLVED: the 35B divergence is a BIT-EXACT tie in the oracle, not a defect + +vLLM was asked for `logprobs` on the same prompt and its own top-2 at the +divergence position, token 7, are the SAME float32 value: + +| pos | top1 | logprob | top2 | logprob | diff | +|---|---|---:|---|---:|---:| +| 6 | `' its'` | -0.011696805246174335 | `' iconic'` | -5.511696815490723 | 5.500000 | +| 7 | `' rich'` | -1.2221027612686157 | `' iconic'` | -1.2221027612686157 | **0.0** | +| 8 | `' history'` | -0.13290393352508545 | `' cultural'` | -2.507904052734375 | 2.375000 | + +`' rich'` is token 8807 and `' iconic'` is token 25438. Their logits are EQUAL in +the oracle's own arithmetic, so vLLM's `torch.argmax` returns the lower index and +picks `' rich'`. Our sampler implements the same rule -- `ArgReduce` in +`cuda_sample.cu` is documented and written as lowest-index-wins and compares the +true global index, so the tie-break is order-independent -- but our logits are +not bit-identical to vLLM's (no two different GEMM/accumulation orders are), so +the equality does not reproduce on our side and 25438 wins. + +**So this is not a forward defect and the 35B is not broken.** It is the exact +case the ratified near-tie gate exists for: at a zero-margin position the outcome +is decided below the precision of either implementation, and demanding +token-exactness there is demanding that two different kernel stacks produce +bit-identical floats. The earlier framing on this page -- first a state-dependent +bug, then a flat token-exactness failure -- is superseded by this measurement. + +Two things remain true and are the actual residue. The 35B gate certifies the +forward with ONE pinned prompt, and the second prompt tried lands on a zero-margin +tie, so the gate's coverage is thin and a prompt set with margins above the noise +floor would be worth more than the single prompt. And the 35B throughput table +stands: ours 70.58 / 107.64 / 194.87 / 318.56 against vLLM's 74.71 / 130.35 / +204.67 / 354.04, so 0.945 / 0.826 / 0.952 / 0.900 -- we are BEHIND on the 35B +too, by 5 to 10 percent on the trustworthy cells, and that is now unblocked +ordinary speed work rather than a correctness question. + +### 35B-A3B decode attribution (2026-08-09) + +Two-length `nsys` difference at `--cuda-graph-trace=node`, same discipline as the +27B. Decode-only GPU kernel time is 16.686 ms/step. + +| Share | Calls/step | us/step | Kernel | +|---:|---:|---:|---| +| 19.8% | 40 | 3308 | `nvjet_sm121_qqsss_mma_128x64x128` (cuBLASLt FP8) | +| 18.6% | 80 | 3097 | `marlin_moe_wna16::Marlin` (routed experts) | +| 10.0% | 40 | 1661 | `nvjet_sm121_qqtst_mma_128x64x128_splitK` | +| 9.7% | 41 | 1611 | `marlin::Marlin` (the DENSE route) | +| 7.2% | 30 | 1198 | `nvjet_sm121_qqsss_mma_192x16x128` | +| 5.6% | 40 | 941 | cuBLAS `gemvx` bf16 | +| 4.7% | 40 | 780 | `MoeRouterTopKKernel` | +| 3.6% | 30 | 603 | `GdnDecodeFusedKernel` | +| 3.1% | 41 | 515 | `CastF32Kernel` | +| 1.6% | 40 | 274 | `RmsNormQuantFp8RowKernel` | +| 1.5% | 80 | 249 | `SiluAndMulKernel` | + +Unlike the 27B, no single line dominates: the two GEMM families are 37% and 28% +and the rest is a long tail. Two entries stand out as pure overhead rather than +work. `CastF32Kernel` is 3.1% for 41 calls, one per NVFP4 marlin GEMM, and exists +only because the caller asks for an f32 output from a kernel whose `c_type` is +bf16 -- the same f32-stream pattern the Laguna campaign identified, and here it +is on a GATE model rather than only on laguna and ds4 as the current +`NOW.md` row claims. `MoeRouterTopKKernel` is 4.7% for routing alone. + +Those two together are 7.8% of the step against a 5 to 10 percent deficit, so +they are the first place to look for 35B parity, ahead of touching either GEMM +family. Both need their own gate: dropping the f32 cast changes the residual +stream dtype and therefore the numerics, so it is a spike with a strict 27B/35B +gate, not a one-liner. + +### LANDED: warp-shuffle MoE router top-k (2026-08-09) + +The 35B attribution put `MoeRouterTopKKernel` at 4.7% of the decode step, 780 +us/step over 40 calls, which is 19.5 us for a few hundred comparisons. The cause +is structural rather than algorithmic: the grid is one block per token, so at +decode ONE block on ONE SM ran the k selection rounds, and each round did a +block-wide tree reduction costing log2(256) `__syncthreads`. For k=8 over 256 +experts that is roughly 64 barriers, and barrier latency was the kernel. + +The k rounds now reduce with `__shfl_down_sync` inside each warp and then one +cross-warp pass, taking barriers per round from log2(kBlock)+1 to 2. + +**This is byte-identical, and the reason it is safe is worth stating.** The top-k +rounds are an ARGMAX reduction, not an arithmetic one: the comparison is +unchanged (higher value wins, exact tie goes to the lower expert index) and +argmax under a total order is associative and commutative, so any reduction order +yields the same (value, index). The softmax max and sum reductions in the same +kernel ARE arithmetic and were deliberately left on their existing tree, because +reassociating those would move the denominator by an ulp and could change which +experts win. + +MEASURED, same harness/checkpoint/box, 3 reps, `nvidia/Qwen3.6-35B-A3B-NVFP4` +@`491c2f1e`: + +| c | before | after | gain | vs vLLM before | after | +|---|---:|---:|---:|---:|---:| +| 1 | 70.58 | 73.24 | +3.8% | 0.945x | **0.980x** | +| 2 | 107.64 | 113.02 | +5.0% | 0.826x | **0.867x** | +| 4 | 194.87 | 199.90 | +2.6% | 0.952x | **0.977x** | +| 8 | 318.56 | 325.36 | +2.1% | 0.900x | **0.919x** | + +The 32-token greedy continuation is unchanged from the pre-spike baseline, as the +byte-identity argument requires. Gates on this build: `test_ops_moe` 33451/33451, +`test_ops_moe_grouped` 440/440, `test_ops_moe_grouped_bf16` 19/19, +`test_qwen36_paged_engine` 315/315, `test_qwen27_paged_engine` 235/235. + +c1 and c4 are now within 2 to 3 percent of vLLM. c2 and c8 remain the weak cells +and both carry the wider run-to-run spread, so the next 35B work should start by +tightening those measurements before chasing them. `CastF32Kernel` at 3.1% +remains unclaimed and is the next lever. + +### Grouped-router top-k parallelised; Kimi-Linear effect NOT ESTABLISHED (2026-08-09) + +`MoeRouterGroupedTopKKernel` is the router Kimi-Linear-48B and the DeepSeek-style +MoEs take (`use_grouped_topk`, so `num_expert_group > 0` routes here rather than +to the ungrouped kernel fixed earlier today). It did `if (threadIdx.x != 0) +return;` before the group scoring, so the group scores, the group mask, the top-k +and the renorm ALL ran on one lane of one block. Its own comment said so and +deferred the fix: "a few thousand serial ops per token ... Speed work belongs to +W9, after the numerics are gated". At Kimi-Linear's shape (e=256, k=8) step (4) +alone is ~2048 serial compares per token. + +Step (4) is now block-parallel with the same warp-shuffle argmax used for the +ungrouped router, on the same byte-identity argument: the comparison is unchanged +(higher score wins, exact tie to the lower expert index, which the serial +ascending scan with strict `>` also produced) and argmax over a total order +reassociates freely. Everything ARITHMETIC still runs on thread 0 in the original +order -- the `denom` accumulation in k order, the renormalize, and the +`routed_scaling_factor` -- and steps (2) and (3) stay serial because they are +O(n_group * group_size) once and n_group is 1 or 8 for the shapes we serve. + +Gates on this build: `test_ops_moe` 33451/33451, `test_ops_moe_grouped` 440/440, +`test_ops_moe_grouped_bf16` 19/19, `test_deepseek_v4_moe` 716/716, +`test_qwen36_paged_engine` 315/315, `test_qwen27_paged_engine` 235/235. + +**The speed effect on Kimi-Linear is NOT ESTABLISHED, and the first run said +otherwise.** Same box, same golden, `kimi-linear-gen` two-length diff, 128 steps: + +| rep | baseline tok/s | spike tok/s | paired | +|---|---:|---:|---:| +| 1 | 17.93 | 18.36 | +2.4% | +| 2 | 18.14 | 17.73 | -2.3% | +| 3 | 16.43 | 18.13 | +10.3% | + +Baseline median 17.93 with a 1.104 spread, spike median 18.13 with 1.036. Two of +three pairs favour the spike and one does not, and the baseline's own range +(16.43 to 18.14) is wider than the effect being measured. Reporting the rep-1 ++2.4% would have been reporting noise. Token ids are IDENTICAL in all six runs +(122/128 against the golden, the same `got:` sequence), which is what the +byte-identity argument requires and is the part that IS established. + +The instrument is the limitation, not the change: `kimi-linear-gen`'s two-length +diff carries a ~10% band here, while the 35B's warm-server harness ran spreads +under 1.05 and resolved a 2-5% effect cleanly. Kimi-Linear cannot use that +harness today because the published checkpoint ships tiktoken +(`tiktoken.model` + `tokenization_kimi.py`) and no `tokenizer.json`, so the +server refuses to load it; the measurement above only works against a +hand-prepared `~/kimi-linear-engine-dir` with a converted tokenizer. Giving +Kimi-Linear a warm-server path is a prerequisite for any binding Kimi speed +number, and is worth more than re-running this A/B. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index c48084c8d..35001ab51 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -14,13 +14,15 @@ | **DSR fix: async readback capability (2026-08-08)** | **No number owed**: behavior-neutral (CPU/CUDA async-ON, discrete non-CUDA async-OFF, unchanged); moves a `kCUDA` check onto `Backend`, unblocking red CI on #127/#154/#155 | | **`ROAD-V1-MEM` M1+M2 (2026-08-08)** | KV auto-sizing CPU brick: `--kv-cache-memory` sizes the pool from a byte budget via the group-aware `KVBytesPerBlock` divisor (ABI v16, CPU-gated). M3 profile run dgx-gated | | **Record/checker repair 2026-08-07–08** | Gates fixed. Public: `VT_GEMMA4_EXPERT_VRAM_MB` (positive-MiB LRU cap; unset/0 unlimited), `VT_SERVER_MAX_{PROMPT_CHARS,NEW_TOKENS}` (200000/4096; 0 disables); nine Gemma4/ROCm tuners internal. No runtime/perf change. | -| **vLLM** | Qwen3.6-27B NVFP4, GB10 | ahead 4.5% at c1, **tie** at c2 to c32 | identical | -| **vLLM** | Qwen3.6-35B-A3B NVFP4, GB10 | 0.93x to 1.03x: ahead at c4, worst c16 0.93x | identical | +| **vLLM** | Qwen3.6-27B NVFP4 `unsloth` @`890bdef7`, GB10 | ahead 4.5% at c1, **tie** at c2 to c32 | identical | +| **vLLM** | Qwen3.6-27B NVFP4 `nvidia` @`0893e160` (ModelOpt `modelopt_mixed`), GB10 | **0.85x, BEHIND** at every concurrency (0.843x to 0.861x), up from 0.72x | identical | +| **vLLM** | Qwen3.6-35B-A3B NVFP4 `nvidia` @`491c2f1e`, GB10 | decode 0.98x at c1 and c4, 0.87x at c2, 0.92x at c8 (warp-shuffle router landed) | near-tie | | **vLLM** | DeepSeek-V2-Lite (MLA), GB10 | 0.86x to 0.95x throughput, TTFT wins at c4/c8 | identical | | **vLLM** | Laguna-S-2.1 NVFP4 (118B/8B MoE), GB10 | **parity+, 1.03x** (44.46 vs 43.10 tok/s, byte-exact, default config; bf16 weights now device-resident) | near-tie | | **llama.cpp** | Qwen3.5-2B GGUF, CPU aarch64 | 20-core Arm/i8mm: prefill **1.18x ahead**, decode tie, memory parity. RPi5/A76: vllm.cpp is **0.461x prefill / 0.653x decode+E2E**, but uses **24.2% less RSS** | byte-identical on both Arm lanes | | **MLX-LM** | Qwen3-0.6B, Apple M4 | 97.6% warm total, prefill ahead | near-tie | | **DwarfStar** | DeepSeek-V4-Flash GGUF, GB10 | **beats ds4, 1.144x** (18.69 vs 16.33 tok/s, byte-exact, default config) | n/a, GGUF peer | +| **vLLM** | Kimi-Linear-48B-A3B, GB10 | no binding number: the published checkpoint is tiktoken-only, so it cannot drive the warm-server harness | golden 122/128, near-tie profile | Reading the ratios: throughput is ours/reference, latency is reference/ours, so **1.0 or higher is a win** everywhere on this page. Which architecture each number @@ -34,7 +36,8 @@ The binding comparison. vLLM runs its **production graphed config**, never | Model | Quant | vLLM pin | Axes passing | Disposition | |---|---|---|---:|---| -| Qwen3.6-27B | NVFP4 | 0.25.0 | **115/124** | Effective parity-or-better, two-grid totality. Measured on `unsloth/Qwen3.6-27B-NVFP4` @`890bdef7` (BF16 head); @`ccdaab7e` re-quantized the head to FP8 | +| Qwen3.6-27B | NVFP4 (`unsloth` @`890bdef7`) | 0.25.0 | **115/124** | Effective parity-or-better, two-grid totality. Revision-PINNED (the gate no longer lets `readdir` choose): @`ccdaab7e` is the same repo name re-quantized to FP8 W8A8 throughout, not NVFP4 | +| Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | | Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks ON: total **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**; local A/B **+2.152%** ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | @@ -51,7 +54,11 @@ The binding comparison. vLLM runs its **production graphed config**, never ### Qwen3.6-27B by concurrency Medians of three interleaved repetitions, 1,024 in / 128 out, cache off, closed -loop. Output is token-for-token identical to vLLM at every point. +loop. Output is token-for-token identical to vLLM at every point. The `nvidia` +ModelOpt table that follows is a DIFFERENT checkpoint on a DIFFERENT axis and +must not be compared against this one: a 7-token prompt scored on OUTPUT tokens +per second, with method, attribution and the owed 1,024 in / 128 out axis in +[the benchmark record](../.agents/benchmark-record.md). | Concurrency | 1 | 2 | 4 | 8 | 16 | 32 | |---|---:|---:|---:|---:|---:|---:| @@ -68,6 +75,21 @@ low-concurrency *median* decode and TTFT, and wins the corresponding *tail* and the same metric at higher concurrency (c8 p99 ITL 0.86x, but 1.055x at c16 and 1.078x at c32). +### Qwen3.6-27B NVFP4 `nvidia` @`0893e160` by concurrency (ModelOpt) + +| Concurrency | 1 | 2 | 4 | 8 | +|---|---:|---:|---:|---:| +| **vllm.cpp** tok/s | 10.41 | 20.22 | 38.76 | 72.49 | +| vLLM 0.25.0 tok/s | 12.29 | 23.49 | 45.42 | 86.01 | +| **Ratio** | **0.847x** | **0.861x** | **0.853x** | **0.843x** | +| Before the FP8 tower fix | 8.76 | 17.07 | 33.01 | 62.13 | +| Spread, ours / vLLM | 1.000 / 1.006 | 1.009 / 1.069 | 1.001 / 1.001 | 1.005 / 1.003 | +| Method | medians of 3, warm servers, one `flock`, greedy, `ignore_eos` so both emit exactly 128 tokens, `--gpu-memory-utilization 0.55 --max-model-len 4096`, vLLM in its production graphed config | | | | +| Tokens | greedy continuation IDENTICAL between engines, captured from the same warm processes as these numbers | | | | +| Reading | still a real loss, not a tie, and still FLAT across the sweep, so batching and the scheduler are not the cause | | | | +| 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 | | | | + ### Qwen3.6-35B-A3B by concurrency | Concurrency | 1 | 2 | 4 | 8 | 16 | 32 | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 1b7ff2610..6aa2f11fb 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -89,6 +89,7 @@ portable/reference path. In normal operation leave them unset. | `VT_GDN_PACKED_DECODE` | on (CUDA GDN) | Unpacked GDN decode path | | `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_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 8dc9453b7..322d0f8a6 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -95,7 +95,7 @@ speed-pending, which [BENCHMARKS.md](BENCHMARKS.md) tracks. | Architecture | Tested checkpoint(s) | Correctness gate | Speed vs reference | |---|---|---|---| -| `Qwen3_5ForConditionalGeneration` | Qwen3.6-27B NVFP4; Qwen3.5-4B BF16 | 27B strict 235/235 text + 32/32 image/video; 4B cached 3/3 | 27B at/above vLLM; 4B throughput 1.021x, latency/VRAM pending. `lm_head` loads BF16, FP8 or NVFP4 (#164). CUDA/CPU only; the off-CUDA host-pointer bug (#125) is fixed but unrun | +| `Qwen3_5ForConditionalGeneration` | Qwen3.6-27B NVFP4 (`unsloth` @`890bdef7`, `nvidia` @`0893e160`); Qwen3.5-4B BF16 | 27B strict 235/235 text + 32/32 image/video; 4B cached 3/3 | `unsloth` 27B at/above vLLM, `nvidia` ModelOpt 0.85x; 4B throughput 1.021x. Loads BF16, FP8 and NVFP4 (CT + ModelOpt naming); a `modelopt_mixed` FP8 tower stays NATIVE (#164). CUDA/CPU only | | `Qwen3_5MoeForConditionalGeneration` | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | strict 315/315 text vs vLLM 0.25.0 | gate model: 0.93x to 1.03x grid | | `Qwen3ForCausalLM` | Qwen3 dense 0.6B/1.7B/4B/32B, NVFP4A16 | near-tie strict 16/16 vs vLLM 0.25.0 | c1 every-axis parity, c8 decode residual | | `Qwen3MoeForCausalLM` | Qwen3-Coder-30B-A3B | strict 6/6 vs vLLM 0.25.0 | 11/16 grid cells at or above graphed vLLM | diff --git a/docs/STATUS.md b/docs/STATUS.md index 8324b55da..81faaad8d 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -75,8 +75,8 @@ token-for-token correctness against the pinned oracle. | Capability | State | Notes | |---|---|---| -| Qwen3.6-27B (NVFP4) text generation | Correctness-complete, at/above vLLM speed | Token-exact greedy GB10; beats vLLM 0.25.0 tput every c (1.007-1.045x), parity 115/124. FP8/NVFP4 heads load (#164); C10 needs a pin move | -| Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; 3-rep grid 0.93-1.03x. Async batch-1 token-0 degeneration FIXED: `VT_ASYNC_DEVICE_MIRROR` default 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.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-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)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | @@ -86,7 +86,7 @@ token-for-token correctness against the pinned oracle. | DeepSeek-V2 MLA | Correctness-complete, speed-pending | Token-exact 8/8 (DeepSeek-V2-Lite); 0.86-0.95x output rate, TTFT faster at c4/c8. A2+A5 MLA norm-rope fold default-ON (`VT_MLA_FUSED_NORM_ROPE`, bit-exact rollback, SACRED 8/8 unchanged; forensics in benchmark-record) — kimi_k3/kimi-linear inherit it | | GLM-4 dense (sandwich norms, partial rope) | Correctness-complete, speed-pending | Token-exact 16/16 (GLM-4-9B-0414); first GLM-family model; partial interleaved RoPE + Gemma2 sandwich norms + biased qkv | | GLM-4.7-Flash (MLA + GLM MoE) | Correctness-complete, speed-pending | Token-exact 8/8 (GLM-4.7-Flash, 31.2B); reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora query branch + noaux_tc sigmoid router with routed-scaling | -| Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE hybrid) | **RUNNER FOLD LANDS (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; golden 122/128 (near-tie profile); FA2 MLA default-ON; `vllm_complete_tokens` (ABI v13).** STRICT stays CLOSED. Server 19.0 tok/s wall (~0.90× vLLM floor) = speed open | paged suite 8/8·206; SACRED post-fold 35B 315/315 + 27B 235/235; thin ABI client (ratchet 8) | +| Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE hybrid) | **RUNNER FOLD LANDS (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; golden 122/128 (near-tie profile); FA2 MLA default-ON; `vllm_complete_tokens` (ABI v13).** Grouped-router top-k block-parallel (byte-identical); no binding speed number: ckpt is tiktoken-only, so no warm-server harness. STRICT stays CLOSED. Server 19.0 tok/s wall (~0.90× vLLM floor) = speed open | paged suite 8/8·206; SACRED post-fold 35B 315/315 + 27B 235/235; thin ABI client (ratchet 8) | | Gemma-3 dense (GeGLU, dual rope, sandwich norms) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-3-1b-it); first Gemma-family model; GeGLU (gelu_pytorch_tanh) + dual per-layer RoPE theta + Gemma-RMSNorm sandwich norms + sqrt(hidden) embed-scale + query_pre_attn_scalar scaling | | Gemma-2 dense (attn + final logit soft-cap) | Correctness-complete, speed-pending | Near-tie-band 48/48 (gemma-2-2b-it): 44/48 strict on vLLM's greedy + 4/48 at 0.0-nat ties in vLLM's own logits; proves the attention + final logit soft-cap primitives (attn_logit_softcapping 50 + final 30); the inverse of Gemma-3 (both soft-caps, no QK-norm) | | Gemma-1 dense (the original Gemma) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-2b); two fused norms/layer, head_dim scale, GeGLU + sqrt(hidden) embed-scale, tied lm_head; no soft-cap/QK-norm/sliding. **D1 (2026-07-31): the whole Gemma family (1/2/3/4) folded to the default-ON bf16 merged-QKV descriptor (`MergedQkvEnabled`); re-gated Gemma-2 SACRED 48/48 (global+sliding) + Gemma-4 STRICT 32/32 — its existing gate held** | @@ -107,7 +107,7 @@ token-for-token correctness against the pinned oracle. | Long-context RoPE + sliding-window attention | Correctness feature-positive on GB10, speed-pending | Shared scaled-RoPE (YaRN, Llama-3, Phi-3/4 LongRoPE, dynamic-NTK) + sliding-window attention, GPU-gated vs the vLLM 0.26 oracle: LongRoPE (Phi-4-mini) + llama3 (Llama-3.2-1B) + dynamic-NTK (InternLM2) 16/16, sliding-window (Gemma-2/Gemma-3) 48/48, operator local-mask kernel positive. YaRN and chunked-local model e2e are reachable-blocked (no cached vehicle); speed pending | | LoRA / multi-LoRA adapters | In progress (W1 CPU brick), not yet usable end-to-end | Highest-demand missing feature. W0 spike + W1 CPU runtime brick landed (`LORA-RUNTIME` ACTIVE): the `LoRALayerWeights` container, the punica shrink/expand ops (`-1`-slot skip), `AddLoraLinear`, and a single-linear `LoRALinear` (ReplicatedLinear create/set/reset/apply, scaling folded at SetLora), RUNTIME-VERIFIED on CPU (`test_punica_cpu` 6/6, 101 assertions, vs double-precision references, RED-first). Packed/TP/merged layers, mapping metadata, adapter load, LRU multi-adapter manager, the load/unload endpoints, and the GPU kernels + model gate are named W2-W7 in `.agents/specs/lora-adapter.md`. No model can be served with an adapter yet | | Safetensors loading | Supported | Both gate models plus every registered dense/MoE family | -| GGUF loading (F32/F16/BF16/Q4_0/Q8_0/Q2_K/Q3_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/IQ3_XXS/IQ2_S/MXFP4/NVFP4) | Supported; compute-in-quant (keep-quant) on CPU AND now CUDA for the six K-block encodings PLUS Q2_K/IQ2_XXS/IQ3_XXS (DeepSeek-V4 W8, 2026-07-29 - the FIRST CUDA keep-quant GGUF k-quant GEMM `KERNEL-QUANT-CIQ-GEMM-CUDA`, MMVQ-style dequant-in-kernel, GB10-gated 92401/92401 vs the CPU oracle, so a CUDA runner keeps blocks compressed and dots them on the GPU instead of the ARM cores); **NVFP4 now COMPUTES IN FP4 on CUDA for the dense-MLP and full-attention projections (2026-07-29, `CLAIM-GGUF-NVFP4-COMPUTE`), no longer materialize-only** | Weights in six block encodings stay compressed from file to matmul on CPU (no BF16 expansion). NVFP4 (ggml type 40) DEQUANTIZES, including the per-tensor (per-expert) `.scale` sidecar the container keeps outside the blocks; gated BIT-EXACT against the compressed-tensors NVFP4 path on real Qwen3.6-27B bytes from both containers. **It no longer expands to bf16 on CUDA:** an NVFP4 matmul/expert weight is REPACKED at load into the same (`weight_packed [N,K/2]`, `weight_scale [N,K/16]`) operand pair the compressed-tensors path produces - a pure byte permutation, gated BYTE-IDENTICAL against that container - and the existing `vt::MatmulNvfp4*` kernels run on it, so no new kernel exists and no numerics are re-derived. Covers the dense MLP + full-attention q/k/v/o and the MoE shared/routed experts; the GDN `in_proj_*` family and `ssm_out` still expand (the V-head reorder rewrites their layout) and a CPU build still expands everything - the documented `part` subset. **MEASURED on GB10 (2026-07-29), same-binary A/B, one `flock`, idle box, 2 reps per arm:** peak RSS **50.8 -> 25.7 GiB**, load-and-generate **1:58 -> 0:41**, and the 256 projections that move cost 35 840 MiB expanded against 10 080 MiB fp4-resident (3.56x). **The divergence against the safetensors sibling CLOSES:** the fp4 arm is token-IDENTICAL over 24 greedy tokens where the bf16 arm of the same binary diverges at index 4, which retires the reading that that divergence was permanent. It is REPORTED, not gated: the two containers are not the same model - the GGUF NVFP4-quantizes 192 GDN `in_proj` tensors the safetensors keeps BF16 (mean relative weight error ~0.18) and their activation global scales differ - so identity is not guaranteed and a cross-container throughput arm is not valid. SACRED gates unmoved: `test_qwen27_paged_engine` 235/235, `test_qwen36_paged_engine` 315/315. **The MoE (35B) stacked-expert arm is now HARDWARE-GATED too (2026-07-29)**, superseding the gap recorded here: the real 35B A3B NVFP4 GGUF loads and generates through the fp4 path, its 120 routed-expert stacks x 256 experts repack to the modelopt safetensors' own operands with ZERO differing bytes over 840 sampled (tensor, expert) slabs, and all 840 per-expert `.scale[e]` are bit-identical to that expert's `weight_scale_2` - the per-expert scale INDEXING, mutation-proved against both a `scales[0]`-for-every-expert and an expert-0-slab-for-every-expert mutant. Same-binary A/B: peak RSS 68.5 -> 22.7 GiB (3.01x), load-and-generate 1:51.9 -> 0:28.8, tokens IDENTICAL (correct here, since the 35B routed experts run the W4A16 grouped GEMM in both arms). Recorded as OPEN, not smoothed over: this case's 24-token greedy stream is NOT run-to-run stable (one of three `use_a16` runs and one of four safetensors-reference runs differed), so the binding results are the weight-level byte identity and the residency audit, not a token-exactness claim; `test_qwen36_paged_engine` is token-exact at ITS engine params, so the instability belongs to this case's configuration and attributing it is owed work. That run also found and FIXED a latent defect the MoE arm made reachable: the two fp4 fused MoE blocks issued the router GEMM assuming the safetensors `[K,N]` gate layout and threw `matmul: inner dims mismatch` on the GGUF's `[N,K]` one; `MoeRouterLogits` now branches on `nk` (inert for the safetensors path, SACRED gates unmoved). **Q2_K (id 10) + IQ2_XXS (id 16) DEQUANTIZE (2026-07-29, `CLAIM-DSV4-GGUF-LOADER`):** the ~2-bit types the single-Spark `DeepSeek-V4-Flash-GGUF UD-IQ2_XXS`/`UD-Q2_K_XL` vehicles use, ported 1:1 from llama.cpp `ggml-quants.c` (`iq2xxs_grid` codebook + signs; Q2_K nibble sub-scale/min), unit-gated on hand-derived bytes (`test_gguf_dequant` 15/15). Dequant-only (no vec_dot -> expand-bf16). A V4-GGUF model still cannot RUN: the V4-GGUF name map (tensor-manifest-blocked) + the V4 forward (W3-W8) remain. **Multi-shard split GGUF READING landed (2026-08-03, `CLAIM-GGUF-SPLIT-SHARDS`):** `GgufFile::Open` now transparently stitches llama.cpp `gguf-split` shards (`...-00001-of-00003.gguf`) — every shard mmap'd, tensor tables merged, KV metadata taken from shard `00001`, and the sibling shard mappings kept alive by the primary so keep-quant mmap-borrows stay valid across shards (`OwnsSpan` is shard-aware); `VT_GGUF_NO_SPLIT=1` opts out; unit-gated (`test_gguf` split-merge / no-split / count-mismatch cases, 33/33 local). This unblocks the real 3-shard `unsloth/DeepSeek-V4-Flash-0731 UD-IQ2_M` (~91 GiB), whose layout is the NATIVE `deepseek4` arch — per-block `ffn_gate_tid2eid` hash tables (hash layers 0/1/2) + `hc_*` MHC + DSA compressor/indexer are all PRESENT (name-map 1328/1328), `vocab_size` derives from `token_embd` — NOT a standard llama.cpp conversion, so no loader-layout change is owed. It now loads THROUGH 1324/1328 tensors; the sole remaining gap is 4 routed-expert slabs quantized with IQ2_S (id 22, ×2) + MXFP4 (id 39, ×2) — encodings we have GGUF block traits for but no keep-quant vec_dot, so they hit the expand→dequant path which lacks them. Dequant-expanding those 4 big expert tensors to bf16 would add ~17 GiB (~106 GiB total → GB10 OOM-reboot risk), so the memory-safe fix is an IQ2_S+MXFP4 keep-quant kernel (CPU dequant dispatch + the `iq2s_grid` codebook + a CUDA `DotSuperblock`), spec'd as the next brick **IQ2_S (id 22) + MXFP4 (id 39) DEQUANTIZE + KEEP-QUANT on CPU (2026-08-03, `CLAIM-DSV4-UDIQ2M-QUANT`, off-GPU):** the extra per-tensor "dynamic" encodings the `unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_M` checkpoint mixes into its last routed-expert slabs (IQ2_S `ffn_gate/up` dotting Q8_K, MXFP4 `ffn_down` dotting Q8_0) — ported 1:1 from llama.cpp `ggml-quants.c` @ 237ad9b96 (`iq2s_grid` 1024-entry codebook + DIRECT sign bytes; MXFP4 `kvalues_mxfp4` + `e8m0_to_fp32_half` micro-scaling, distinct from the compressed-tensors `E8M0ToF32` NVFP4 path). CPU dequant + keep-quant `vec_dot`, unit-gated on hand-derived golden bytes (`test_gguf_dequant` 17/17), an INDEPENDENT f64 dequant-then-dot + GEMM NMSE (`test_ops_quant_dot` 19/19), and keep-quant routing (`test_gguf_keep_quant` 37/37) — all CPU-green, so UD-IQ2_M's four previously-`unsupported ggml type 22/39` slabs now load COMPRESSED (no ~17 GiB bf16 expansion that OOM-reboots the box). CUDA: the IQ2_S device `DotSuperblock` is wired into the Q8_K grouped-MoE GEMM and now **CUDA-BUILT + LINKED on GB10 (sm_121a, CUDA 13.0, `-Werror`, 2026-08-03 integration)** — it compiles clean and the merged binary links; MXFP4's device dot (`DotMXFP4`) is written but NOT wired (Q8_0-activation needs a separate 32-block GEMM) so it is marked `[[maybe_unused]]` to keep the ready math without tripping nvcc #177-D, and on GPU MXFP4 CPU-fallbacks like Q4_0/Q8_0. The V4-GGUF forward + a real UD-IQ2_M GPU load/coherence run are owed | +| GGUF loading (F32/F16/BF16/Q4_0/Q8_0/Q2_K/Q3_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/IQ3_XXS/IQ2_S/MXFP4/NVFP4) | Supported; compute-in-quant (keep-quant) on CPU AND now CUDA for the six K-block encodings PLUS Q2_K/IQ2_XXS/IQ3_XXS (DeepSeek-V4 W8, 2026-07-29 - the FIRST CUDA keep-quant GGUF k-quant GEMM `KERNEL-QUANT-CIQ-GEMM-CUDA`, MMVQ-style dequant-in-kernel, GB10-gated 92401/92401 vs the CPU oracle, so a CUDA runner keeps blocks compressed and dots them on the GPU instead of the ARM cores); **NVFP4 now COMPUTES IN FP4 on CUDA for the dense-MLP and full-attention projections (2026-07-29, `CLAIM-GGUF-NVFP4-COMPUTE`), no longer materialize-only** | Weights in six block encodings stay compressed from file to matmul on CPU (no BF16 expansion). NVFP4 (ggml type 40) DEQUANTIZES, including the per-tensor (per-expert) `.scale` sidecar the container keeps outside the blocks; gated BIT-EXACT against the compressed-tensors NVFP4 path on real Qwen3.6-27B bytes from both containers. **It no longer expands to bf16 on CUDA:** an NVFP4 matmul/expert weight is REPACKED at load into the same (`weight_packed [N,K/2]`, `weight_scale [N,K/16]`) operand pair the compressed-tensors path produces - a pure byte permutation, gated BYTE-IDENTICAL against that container - and the existing `vt::MatmulNvfp4*` kernels run on it, so no new kernel exists and no numerics are re-derived. Covers the dense MLP + full-attention q/k/v/o and the MoE shared/routed experts; the GDN `in_proj_*` family and `ssm_out` still expand (the V-head reorder rewrites their layout) and a CPU build still expands everything - the documented `part` subset. **MEASURED GB10 (2026-07-29), same-binary A/B, 2 reps/arm:** peak RSS **50.8 -> 25.7 GiB**, load-and-generate **1:58 -> 0:41**; the 256 projections that move cost 35 840 MiB expanded vs 10 080 MiB fp4-resident (3.56x). **The safetensors-sibling divergence CLOSES:** the fp4 arm is token-IDENTICAL over 24 greedy tokens where the same binary's bf16 arm diverges at index 4. It is REPORTED, not gated: the two containers are not the same model - the GGUF NVFP4-quantizes 192 GDN `in_proj` tensors the safetensors keeps BF16 (mean relative weight error ~0.18) and their activation global scales differ - so identity is not guaranteed and a cross-container throughput arm is not valid. SACRED gates unmoved: `test_qwen27_paged_engine` 235/235, `test_qwen36_paged_engine` 315/315. **The MoE (35B) stacked-expert arm is now HARDWARE-GATED too (2026-07-29)**, superseding the gap recorded here: the real 35B A3B NVFP4 GGUF loads and generates through the fp4 path, its 120 routed-expert stacks x 256 experts repack to the modelopt safetensors' own operands with ZERO differing bytes over 840 sampled (tensor, expert) slabs, and all 840 per-expert `.scale[e]` are bit-identical to that expert's `weight_scale_2` - the per-expert scale INDEXING, mutation-proved against both a `scales[0]`-for-every-expert and an expert-0-slab-for-every-expert mutant. Same-binary A/B: peak RSS 68.5 -> 22.7 GiB (3.01x), load-and-generate 1:51.9 -> 0:28.8, tokens IDENTICAL (the 35B routed experts run the W4A16 grouped GEMM in both arms). Recorded OPEN: this case's 24-token greedy stream is NOT run-to-run stable (1 of 3 `use_a16` and 1 of 4 reference runs differed), so the binding results are the weight-level byte identity and the residency audit, not token-exactness; `test_qwen36_paged_engine` is token-exact at ITS engine params, so the instability belongs to this case's configuration and attributing it is owed work. It also FIXED a latent defect the MoE arm made reachable: the two fp4 fused MoE blocks issued the router GEMM assuming the safetensors `[K,N]` gate layout and threw `matmul: inner dims mismatch` on the GGUF's `[N,K]`; `MoeRouterLogits` now branches on `nk` (inert for safetensors, SACRED unmoved). **Q2_K (id 10) + IQ2_XXS (id 16) DEQUANTIZE (2026-07-29, `CLAIM-DSV4-GGUF-LOADER`):** the ~2-bit types the single-Spark `DeepSeek-V4-Flash-GGUF UD-IQ2_XXS`/`UD-Q2_K_XL` vehicles use, ported 1:1 from llama.cpp `ggml-quants.c` (`iq2xxs_grid` codebook + signs; Q2_K nibble sub-scale/min), unit-gated on hand-derived bytes (`test_gguf_dequant` 15/15). Dequant-only (no vec_dot -> expand-bf16). A V4-GGUF model still cannot RUN: the name map (tensor-manifest-blocked) + the V4 forward (W3-W8) remain. **Multi-shard split GGUF READING landed (2026-08-03, `CLAIM-GGUF-SPLIT-SHARDS`):** `GgufFile::Open` now transparently stitches llama.cpp `gguf-split` shards (`...-00001-of-00003.gguf`) — every shard mmap'd, tensor tables merged, KV metadata taken from shard `00001`, and the sibling shard mappings kept alive by the primary so keep-quant mmap-borrows stay valid across shards (`OwnsSpan` is shard-aware); `VT_GGUF_NO_SPLIT=1` opts out; unit-gated (`test_gguf` split-merge / no-split / count-mismatch cases, 33/33 local). This unblocks the real 3-shard `unsloth/DeepSeek-V4-Flash-0731 UD-IQ2_M` (~91 GiB), whose layout is the NATIVE `deepseek4` arch — per-block `ffn_gate_tid2eid` hash tables (hash layers 0/1/2) + `hc_*` MHC + DSA compressor/indexer are all PRESENT (name-map 1328/1328), `vocab_size` derives from `token_embd` — NOT a standard llama.cpp conversion, so no loader-layout change is owed. It now loads THROUGH 1324/1328 tensors; the sole remaining gap is 4 routed-expert slabs quantized with IQ2_S (id 22, ×2) + MXFP4 (id 39, ×2) — encodings we have GGUF block traits for but no keep-quant vec_dot, so they hit the expand→dequant path which lacks them. Expanding those 4 expert tensors to bf16 would add ~17 GiB (~106 GiB total → OOM-reboot risk), so the memory-safe fix is an IQ2_S+MXFP4 keep-quant kernel (CPU dequant dispatch + `iq2s_grid` + a CUDA `DotSuperblock`), spec'd as the next brick **IQ2_S (id 22) + MXFP4 (id 39) DEQUANTIZE + KEEP-QUANT on CPU (2026-08-03, `CLAIM-DSV4-UDIQ2M-QUANT`, off-GPU):** the extra per-tensor "dynamic" encodings the `unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_M` checkpoint mixes into its last routed-expert slabs (IQ2_S `ffn_gate/up` dotting Q8_K, MXFP4 `ffn_down` dotting Q8_0) — ported 1:1 from llama.cpp `ggml-quants.c` @ 237ad9b96 (`iq2s_grid` 1024-entry codebook + DIRECT sign bytes; MXFP4 `kvalues_mxfp4` + `e8m0_to_fp32_half` micro-scaling, distinct from the compressed-tensors `E8M0ToF32` NVFP4 path). CPU dequant + keep-quant `vec_dot`, unit-gated on hand-derived golden bytes (`test_gguf_dequant` 17/17), an INDEPENDENT f64 dequant-then-dot + GEMM NMSE (`test_ops_quant_dot` 19/19), and keep-quant routing (`test_gguf_keep_quant` 37/37) — all CPU-green, so UD-IQ2_M's four previously-`unsupported ggml type 22/39` slabs now load COMPRESSED (no ~17 GiB bf16 expansion that OOM-reboots the box). CUDA: the IQ2_S device `DotSuperblock` is wired into the Q8_K grouped-MoE GEMM and now **CUDA-BUILT + LINKED on GB10 (sm_121a, CUDA 13.0, `-Werror`, 2026-08-03 integration)** — it compiles clean and the merged binary links; MXFP4's device dot (`DotMXFP4`) is written but NOT wired (Q8_0-activation needs a separate 32-block GEMM) so it is marked `[[maybe_unused]]` to keep the ready math without tripping nvcc #177-D, and on GPU MXFP4 CPU-fallbacks like Q4_0/Q8_0. The V4-GGUF forward + a real UD-IQ2_M GPU load/coherence run are owed | | AWQ / GPTQ quantization | W0 spike + W1 CPU INT4 dequant primitive; not yet loadable end to end | INT4 unpack+dequant-to-bf16 for BOTH community formats, mirroring vLLM 1:1 (AWQ reverse-order `awq_triton.py`; GPTQ `qdq_4.cuh` with zero_offset v1/v2 + act-order g_idx). Unit-gated RED-first (hand-computed known bytes + double-precision roundtrip). NOT wired to a loader, no GPU Marlin compute, no model run yet: config recognizer (W2), Marlin GPU GEMM riding the vendored NVFP4 Marlin (W4), CPU e2e (W3), GPTQ 8/2/3-bit (W5) and MoE (W6) are named next bricks. See [.agents/specs/awq-gptq-quant.md](../.agents/specs/awq-gptq-quant.md) | | MXFP4 (compressed-tensors `mxfp4-pack-quantized`) | Compute PROVEN (#38); GQA-swap ON (#49); decode-graph+gate_up FUSION default-ON. `VT_MARLIN_DENSE` DEFAULT-ON (`KERNEL-MARLIN-DENSE-EXEC`): dense marlin 48-CTA byte-faithful (32B 0.000, 263/263), binding beats #51 every axis (c1 1.020, c8 0.969, mem 2.63x). **`QUANT-CT-MXFP4-FINAL-STACK` TERMINAL — both last levers exhausted: num_splits cap `VT_FA2_NSPLITS_CAP` gated-OFF (c1-only, self-corrects@c8; 32B strict char-identical); glue folds via `vt::FusedChain`; `FLASH-AUDIT` #68: c8 flash gap +12.5us/call is occupancy/L2-bound; `-use_fast_math` TRIED, REGRESSES flash (168.8→189.8), rejected. c1 1.020x PASS, c2-c8 0.962-0.969.** state.md | Shared with DeepSeek-V4-Flash + Kimi-K3 MXFP4 paths. CPU E8M0 dequant 5/5·1142. GPU W4A4 + MoE-expert e2e later | | CPU backend vs llama.cpp | 20-core at floor; RPi5/A76 below floor `GATING` | Pi: AAPCS64 beats SDOT 3.66-5.08%; llama.cpp 2.17x pf / 1.53x dec faster (0.461x/0.653x); RSS -24.2%; 64-tok byte-exact. BF16 GEMM open | diff --git a/docs/USAGE.md b/docs/USAGE.md index 06a2a0e00..a5ec7007f 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -78,10 +78,14 @@ a silent fallback cannot post a plausible number: reports GPU-timestamp time rather than wall clock; see [ENVIRONMENT.md](ENVIRONMENT.md) for what each knob does and what it measured. -### Quantized checkpoints: which `lm_head` forms load - -Publishers do not agree on how the output head is stored, and a single repo can -change it between revisions. For the Qwen3.6 dense family we accept all three +### Quantized checkpoints: which weight forms load + +Publishers do not agree on how weights are stored, and a single repo can change +it between revisions (one 27B "NVFP4" repo silently became FP8 throughout). +The table below is about `lm_head`; the same three forms are accepted for the +attention, MLP and `linear_attn` projections, in both compressed-tensors +(`weight_packed` + `weight_global_scale`) and ModelOpt (`weight` + +`weight_scale_2`) naming. For the Qwen3.6 dense family we accept all three forms in use, so pick a checkpoint by its quality, not by its head: | `lm_head.weight` | Companion tensors | Seen in | @@ -271,7 +275,7 @@ to one built without video support. See | `--block-size N` | `32` | KV block size | | `--num-blocks N` | `256` | KV blocks | | `--max-model-len N` | `0` (config default) | Max sequence length | -| `--max-num-seqs N` | `8` | Max concurrent sequences (also sizes the HTTP worker pool) | +| `--max-num-seqs N` | `32` | Max concurrent sequences (also sizes the HTTP worker pool). Was `8`, which put a c8 client exactly on the batch ceiling; vLLM's own default is 1024, which we do not mirror because this also caps the padded decode-graph set | | `--max-num-batched-tokens N` | `0` (per-arch default) | Per-step token budget | | `--enable-prefix-caching` / `--no-enable-prefix-caching` | model default | Override automatic prefix caching | | `--scheduling-policy fcfs\|priority\|lpm` | `fcfs` | Scheduler policy (`lpm` is the SGLang cache-aware policy, see [docs/SGLANG-COMPAT.md](SGLANG-COMPAT.md)) | diff --git a/include/vllm/entrypoints/model_loader.h b/include/vllm/entrypoints/model_loader.h index f239ef78c..494bc9fd1 100644 --- a/include/vllm/entrypoints/model_loader.h +++ b/include/vllm/entrypoints/model_loader.h @@ -81,7 +81,15 @@ struct EngineParams { // CacheConfig.kv_cache_memory_bytes (cache.py:182,189). int64_t kv_cache_memory_bytes = 0; int max_model_len = 0; // 0 => config.max_position_embeddings. - int max_num_seqs = 8; // max concurrent sequences. + // max concurrent sequences. vLLM's default is 1024 (EngineArgs.max_num_seqs); + // ours was 8, which put c8 EXACTLY on the batch ceiling so the 8th stream + // could not co-batch -- measured as a throughput ratio that stayed flat to c4 + // then collapsed at c8. Raising to 32 recovers it (c8 51.66 -> 64.41 tok/s, + // +24.7%, ratio vs vLLM flat ~0.74x instead of degrading to 0.59x). + // NOT vLLM's 1024: max_num_seqs scales KV demand, and GB10's unified memory + // has a narrow usable band (see gb10 OOM/thrash notes). 32 is the value + // MEASURED clean at --gpu-memory-utilization 0.55; higher is unverified. + int max_num_seqs = 32; // Per-step token budget (the chunked-prefill knob). 0 => the bounded PER-ARCH // default (see LoadedEngine::ResolveMaxNumBatchedTokens): dense arch 2048 flat // (vLLM's DEFAULT_MAX_NUM_BATCHED_TOKENS, vllm/config/scheduler.py:42 @ diff --git a/include/vllm/model_executor/models/dense_weight_loaders.h b/include/vllm/model_executor/models/dense_weight_loaders.h index e5040fdf5..c301f3751 100644 --- a/include/vllm/model_executor/models/dense_weight_loaders.h +++ b/include/vllm/model_executor/models/dense_weight_loaders.h @@ -28,6 +28,7 @@ #include #include "vllm/model_executor/model_loader/safetensors_reader.h" // StTensor, MaybeReleaseSourcePages +#include "vllm/model_executor/model_loader/nvfp4_dequant.h" #include "vllm/model_executor/models/qwen3_5_weights.h" // OwnedTensor, TensorResolver #include "vllm/model_executor/models/tensor_parallel.h" // TensorParallel/TpShard (W2) #include "vt/dtype.h" @@ -59,17 +60,73 @@ inline void TransposeBf16(const uint16_t* src, int64_t rows, int64_t cols, } } +// --- FP8 shard materialization (shared by every BF16 loader below) ----------- +// The 2026-08 Qwen3.6-27B NVFP4 republishes quantize parts of the tower to +// per-tensor or per-output-channel FP8 while leaving the rest BF16, and they do +// not agree on WHICH parts: nvidia/Qwen3.6-27B-NVFP4 ships FP8 `linear_attn` +// in_proj_qkv/in_proj_z/out_proj with scalar F32 scales next to NVFP4 attention +// and MLP, while unsloth @ccdaab7e went FP8 across the whole tower with BF16 +// per-output-channel scales. Rather than teach each loader its own dtype rules, +// every BF16 entry point routes its source bytes through this one materializer. +// +// Returns a pointer to BF16 [rows, cols] bytes: the mmap'd source itself when the +// tensor is already BF16 (zero copy, unchanged behavior), or `staging` after +// dequantization. A per-output-channel scale read as per-tensor would be +// silently WRONG rather than loud, so the element count decides and anything +// else is rejected. +inline const uint8_t* MaterializeBf16Source(const TensorResolver& get, + const std::string& name, + const StTensor& t, + std::vector* staging) { + if (t.dtype == "BF16") return static_cast(t.data); + VT_CHECK(t.dtype == "F8_E4M3", + "dense loader: unsupported dtype '" + t.dtype + "' for " + name + + "; supported: BF16, F8_E4M3 (+ _scale)"); + VT_CHECK(t.shape.size() == 2, + "dense loader: expected 2-D weight for FP8 " + name); + const int64_t rows = t.shape[0]; + const int64_t cols = t.shape[1]; + const StTensor& sc = get(name + "_scale"); + const int64_t n_scale = + static_cast(sc.nbytes) / (sc.dtype == "BF16" ? 2 : 4); + VT_CHECK(n_scale == 1 || n_scale == rows, + "dense loader: " + name + + "_scale must be per-tensor or one value per output row"); + staging->resize(static_cast(rows) * static_cast(cols)); + for (int64_t r = 0; r < rows; ++r) { + const int64_t si = (n_scale == 1) ? 0 : r; + float scale = 1.0F; + if (sc.dtype == "BF16") { + uint16_t h = 0; + std::memcpy(&h, static_cast(sc.data) + si * 2, 2); + const uint32_t bits = static_cast(h) << 16; + std::memcpy(&scale, &bits, sizeof(scale)); + } else { + std::memcpy(&scale, static_cast(sc.data) + si * 4, + sizeof(scale)); + } + DequantFp8ToBf16(static_cast(t.data) + r * cols, scale, cols, + staging->data() + static_cast(r) * cols); + } + // Deliberately NOT releasing here: every caller already calls + // MaybeReleaseSourcePages(t.data, t.nbytes) exactly once on the same range. + return reinterpret_cast(staging->data()); +} + // BF16 tensor copied verbatim (optionally reshaped). inline OwnedTensor LoadBf16Direct(const TensorResolver& get, const std::string& name, const std::vector& shape_override = {}) { const StTensor& t = get(name); - VT_CHECK(t.dtype == "BF16", "dense loader: expected BF16 for " + name); + std::vector staging; + const uint8_t* src = MaterializeBf16Source(get, name, t, &staging); std::vector shape = shape_override.empty() ? t.shape : shape_override; OwnedTensor o = MakeOwned(vt::DType::kBF16, shape); - VT_CHECK(t.nbytes == o.bytes.size(), + const size_t src_bytes = + staging.empty() ? t.nbytes : staging.size() * sizeof(uint16_t); + VT_CHECK(src_bytes == o.bytes.size(), "dense loader: byte-size mismatch for " + name); - std::memcpy(o.bytes.data(), t.data, t.nbytes); + std::memcpy(o.bytes.data(), src, src_bytes); // LOAD-SAFETENSORS: source range now copied-then-dead; drop its resident pages // so the owned mirror never double-resides with the mmap (spec §page-lifetime). MaybeReleaseSourcePages(t.data, t.nbytes); @@ -80,12 +137,13 @@ inline OwnedTensor LoadBf16Direct(const TensorResolver& get, inline OwnedTensor LoadBf16Transposed(const TensorResolver& get, const std::string& name) { const StTensor& t = get(name); - VT_CHECK(t.dtype == "BF16", "dense loader: expected BF16 for " + name); VT_CHECK(t.shape.size() == 2, "dense loader: expected 2-D weight for " + name); + std::vector staging; + const uint8_t* src = MaterializeBf16Source(get, name, t, &staging); const int64_t out_dim = t.shape[0]; const int64_t in_dim = t.shape[1]; OwnedTensor o = MakeOwned(vt::DType::kBF16, {in_dim, out_dim}); - TransposeBf16(reinterpret_cast(t.data), out_dim, in_dim, + TransposeBf16(reinterpret_cast(src), out_dim, in_dim, reinterpret_cast(o.bytes.data())); MaybeReleaseSourcePages(t.data, t.nbytes); return o; @@ -106,7 +164,14 @@ inline OwnedTensor LoadMergedBf16RawNK(const TensorResolver& get, shards.reserve(names.size()); for (const std::string& name : names) { const StTensor& tensor = get(name); - VT_CHECK(tensor.dtype == "BF16", "dense loader: expected BF16 for " + name); + // Both new Qwen3.6-27B NVFP4 publishers quantize the GDN in-projections to + // per-tensor FP8 while leaving in_proj_a/b BF16 (nvidia/Qwen3.6-27B-NVFP4: + // in_proj_qkv F8_E4M3 [10240,5120] + scalar weight_scale/input_scale). The + // shard is materialized to BF16 below so the merge, the TP row split and the + // nk=true MatmulBT orientation stay exactly as they were for a BF16 shard. + VT_CHECK(tensor.dtype == "BF16" || tensor.dtype == "F8_E4M3", + "dense loader: unsupported dtype '" + tensor.dtype + "' for " + name + + "; supported: BF16, F8_E4M3 (+ _scale)"); VT_CHECK(tensor.shape.size() == 2, "dense loader: expected 2-D weight for " + name); VT_CHECK(tensor.shape[0] > 0 && tensor.shape[1] > 0, @@ -139,6 +204,50 @@ inline OwnedTensor LoadMergedBf16RawNK(const TensorResolver& get, sharded_out_dim += r.size(); } + // Materialize any FP8 shard to BF16 before the merge. The scale is either a + // single per-tensor F32 scalar (nvidia/Qwen3.6-27B-NVFP4 in_proj_qkv) or one + // value per output row (unsloth @ccdaab7e, stored BF16). Reading a per-channel + // scale as per-tensor would be silently WRONG rather than loud, so the row + // count decides and anything else is rejected. + std::vector> staged(shards.size()); + std::vector src_ptr(shards.size()); + std::vector src_bytes(shards.size()); + for (size_t i = 0; i < shards.size(); ++i) { + const StTensor& shard = *shards[i]; + if (shard.dtype == "BF16") { + src_ptr[i] = static_cast(shard.data); + src_bytes[i] = shard.nbytes; + continue; + } + const int64_t rows = shard.shape[0]; + const int64_t cols = shard.shape[1]; + const StTensor& sc = get(names[i] + "_scale"); + const int64_t n_scale = + static_cast(sc.nbytes) / (sc.dtype == "BF16" ? 2 : 4); + VT_CHECK(n_scale == 1 || n_scale == rows, + "dense loader: " + names[i] + "_scale must be per-tensor or one " + "value per output row"); + staged[i].resize(static_cast(rows) * static_cast(cols)); + for (int64_t r = 0; r < rows; ++r) { + const int64_t si = (n_scale == 1) ? 0 : r; + float scale = 1.0F; + if (sc.dtype == "BF16") { + uint16_t h = 0; + std::memcpy(&h, static_cast(sc.data) + si * 2, 2); + const uint32_t bits = static_cast(h) << 16; + std::memcpy(&scale, &bits, sizeof(scale)); + } else { + std::memcpy(&scale, static_cast(sc.data) + si * 4, + sizeof(scale)); + } + DequantFp8ToBf16(static_cast(shard.data) + r * cols, scale, + cols, staged[i].data() + static_cast(r) * cols); + } + MaybeReleaseSourcePages(shard.data, shard.nbytes); + src_ptr[i] = reinterpret_cast(staged[i].data()); + src_bytes[i] = staged[i].size() * sizeof(uint16_t); + } + VT_CHECK(sharded_out_dim <= std::numeric_limits::max() / in_dim, "dense loader: merged BF16 element count overflow"); const auto elements = @@ -151,14 +260,13 @@ inline OwnedTensor LoadMergedBf16RawNK(const TensorResolver& get, for (size_t i = 0; i < shards.size(); ++i) { const StTensor& shard = *shards[i]; const size_t full = static_cast(shard.shape[0]) * row_bytes; - VT_CHECK(shard.nbytes == full, + VT_CHECK(src_bytes[i] == full, "dense loader: byte-size mismatch for " + names[i]); const ShardRange& r = ranges[i]; const size_t src_off = static_cast(r.begin) * row_bytes; const size_t copied = static_cast(r.size()) * row_bytes; - std::memcpy(merged.bytes.data() + offset, - static_cast(shard.data) + src_off, copied); - MaybeReleaseSourcePages(shard.data, full); + std::memcpy(merged.bytes.data() + offset, src_ptr[i] + src_off, copied); + if (shard.dtype == "BF16") MaybeReleaseSourcePages(shard.data, full); offset += copied; } VT_CHECK(offset == merged.bytes.size(), diff --git a/include/vllm/model_executor/models/qwen3_5_dense.h b/include/vllm/model_executor/models/qwen3_5_dense.h index f470bf9dd..1aada1c09 100644 --- a/include/vllm/model_executor/models/qwen3_5_dense.h +++ b/include/vllm/model_executor/models/qwen3_5_dense.h @@ -135,6 +135,8 @@ OwnedTensor LoadLmHeadAnyDtype(const TensorResolver& get, const std::function& has, const std::string& name); +Fp8Weight LoadFp8RawShared(const TensorResolver& get, const std::string& proj); + OwnedTensor MaterializeCtNvfp4Bf16Transposed(const TensorResolver& get, const std::string& proj); diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 7256e3152..5f4557365 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -441,7 +441,13 @@ def features_errors(text: str) -> list[str]: # state migration so the reduction cannot become untracked growth headroom. # 243632 since 2026-08-09 (measured 243632): the W1 candidate records local # gencode gates while keeping the real ten-SM archive audit pending. - "chars": 243632, + # + # 243600 since 2026-08-09 (measured 243600): the NVFP4 re-verification + # rewrites the 27B, 35B and Kimi-Linear rows onto their pinned revisions and + # their re-measured ratios, and drops the superseded narrative those rows + # carried. Net -32 against the rebased page, re-pinned byte-tight so the + # reduction cannot become untracked growth headroom. + "chars": 243600, "h2_sections": 11, "long_paragraphs": 82, "oversized_cells": 44, diff --git a/src/capi/vllm_c.cpp b/src/capi/vllm_c.cpp index e0cdae520..378de19b7 100644 --- a/src/capi/vllm_c.cpp +++ b/src/capi/vllm_c.cpp @@ -488,7 +488,7 @@ VLLM_API vllm_model_params vllm_model_params_default(void) { p.block_size = 32; p.num_blocks = 0; // 0 => auto: sized by the v16 knobs, else 256. p.max_model_len = 0; - p.max_num_seqs = 8; + p.max_num_seqs = 32; // see EngineParams::max_num_seqs. p.tool_parser = nullptr; // AUTO-detect from the chat template (ABI v4). p.reasoning_parser = nullptr; // AUTO-detect / disabled (ABI v5). p.speculative_config = nullptr; // speculation disabled (ABI v6). diff --git a/src/vllm/entrypoints/openai/server_main.cpp b/src/vllm/entrypoints/openai/server_main.cpp index e615d39f1..1ac634b2b 100644 --- a/src/vllm/entrypoints/openai/server_main.cpp +++ b/src/vllm/entrypoints/openai/server_main.cpp @@ -147,7 +147,7 @@ struct Args { double gpu_memory_utilization = 0.92; long long kv_cache_memory_bytes = 0; int max_model_len = 0; // 0 => config.max_position_embeddings - int max_num_seqs = 8; + int max_num_seqs = 32; // see model_loader.h: 8 clamped c8 batching. int max_num_batched_tokens = 0; // 0 => per-architecture default. // --device: explicit device selection for the TEXT engine (ARCH-ONE-SURFACE // ROW 8), the vLLM DeviceConfig.device names this build serves: "auto" diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 31d5790da..fc34c1c21 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -35,6 +35,7 @@ #include #include +#include "vllm/model_executor/models/dense_nvfp4_gemm.h" // dense_nvfp4::MarlinDenseEnabled #include "vllm/model_executor/model_loader/nvfp4_dequant.h" #include "vt/backend.h" #ifdef VT_BENCH_PROFILE_CONTROL @@ -2362,11 +2363,37 @@ DBuf MatmulNvfp4MarlinD(Dev d, const Tensor& x, const Nvfp4Weight& w, DType out_ const int64_t M = x.shape[0], K = x.shape[1], N = w.n; MarlinDenseResident& mr = MarlinDenseResidentFor(&w); if (!mr.ready) BuildMarlinDenseResident(d, w, mr); - DenseAlignCache& ac = DenseAlignFor(d, static_cast(M)); int sms = 0; void* ws = DenseMarlinWorkspace(d, &sms); d.b.Memset(d.q, ws, 0, static_cast(sms) * 4 * sizeof(int32_t)); + // VT_MARLIN_DENSE (default ON): the framework-wide dense NVFP4 route this TU + // was forked off. E=1 dense projections go through vLLM's OWN dense marlin + // GEMM instead of the single-expert MoE-marlin route. Same resident + // (mr.w/mr.s/mr.g) and workspace; rank-2 operand views (the dense launcher + // wants [K/16, N*2] / [K/gs, N], not the MoE rank-3 [1, ...]) and direct-A, + // so no moe_align cache. Mirrors dense_nvfp4_gemm.h's MatmulNvfp4MarlinD, + // which qwen3 / olmo2 / deepseek_v2 / minimax_h3 already take. + if (dense_nvfp4::MarlinDenseEnabled() && + vt::OpRegistered(vt::OpId::kMarlinDenseGemm, d.q.device.type)) { + DBuf outd(d, DType::kBF16, {M, N}); + Tensor wqd = MakeTensor(mr.w, DType::kI32, d.q.device, {K / 16, N * 2}); + Tensor scd = MakeTensor(mr.s, DType::kI8, d.q.device, {K / 16, N}); + Tensor ggd = MakeTensor(mr.g, DType::kF32, d.q.device, {1}); + Tensor wstd = MakeTensor(ws, DType::kI32, d.q.device, {sms * 4}); + vt::MarlinDenseArgs dargs{static_cast(M), static_cast(N), + static_cast(K)}; + dargs.group_size = 16; + dargs.mxfp4 = false; + vt::MarlinDenseGemm(d.q, outd.t(), x, wqd, scd, ggd, wstd, dargs); + if (out_dtype == DType::kBF16) return outd; + DBuf outf(d, DType::kF32, {M, N}); + vt::CastF32(d.q, outf.t(), outd.t()); + return outf; + } + + DenseAlignCache& ac = DenseAlignFor(d, static_cast(M)); + // Marlin's output is bf16 (c_type=kBFloat16); an f32 result is the bf16 output // upcast (same value it rounds to — mirror of the cutlass f32-scratch cast). DBuf outbf(d, DType::kBF16, {M, N}); @@ -3727,6 +3754,11 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, MergedGdnBaEnabled(d), indt == DType::kBF16 && outdt == DType::kBF16 && MergedGdnBaOutputDType(true) == DType::kBF16 && + // An FP8 GDN tower feeds fp8 projections that vt::GdnPackedDecode + // rejects ("mixed_qkv/a/b/out must share FP16/BF16/F32 dtype"). + // The unpacked decode -- what the 35B fp8 path already runs -- + // handles them, so an fp8 tower is NOT packed-decode eligible. + w.in_proj_qkv_fp8.Empty() && w.in_proj_z_fp8.Empty() && (state.ssm_state.dtype == DType::kF32 || state.ssm_state.dtype == DType::kF16 || state.ssm_state.dtype == DType::kBF16), 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 ad396d6c4..662199949 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp @@ -314,6 +314,74 @@ OwnedTensor LoadLmHeadAnyDtype(const TensorResolver& get, const TensorExists& ha namespace { +// compressed-tensors and ModelOpt both ship NVFP4, with different names AND a +// different global-scale convention: +// +// compressed-tensors .weight_packed U8 + .weight_scale F8 +// + .weight_global_scale F32 (a DIVISOR) +// ModelOpt .weight U8 + .weight_scale F8 +// + .weight_scale_2 F32 (the SCALE itself) +// +// nvidia/Qwen3.6-27B-NVFP4 is ModelOpt, so every `has(.weight_packed)` +// probe missed it and the whole tower fell through to the BF16 path and died at +// the first U8 tensor. LoadCtNvfp4Raw already reciprocates internally, so the +// ModelOpt scale is passed through as 1/weight_scale_2 to land on the same math +// (identical to the lm_head conversion in LoadLmHeadAnyDtype). +bool IsNvfp4Projection(const TensorExists& has, const std::string& proj) { + return has(proj + ".weight_packed") || has(proj + ".weight_scale_2"); +} + +Nvfp4Weight LoadNvfp4AnyNaming(const TensorResolver& get, const TensorExists& has, + const std::string& proj) { + if (has(proj + ".weight_packed")) return LoadCtNvfp4Raw(get, proj); + + const StTensor& packed = get(proj + ".weight"); + VT_CHECK(packed.dtype == "U8", + "qwen3_5 dense: expected U8 ModelOpt weight for " + proj); + VT_CHECK(packed.shape.size() == 2, + "qwen3_5 dense: expected 2-D ModelOpt weight for " + proj); + const int64_t out_dim = packed.shape[0]; + const int64_t in_dim = packed.shape[1] * 2; + VT_CHECK(in_dim % 16 == 0, + "qwen3_5 dense: NVFP4 in_dim must be a multiple of 16 for " + proj); + const StTensor& ws = get(proj + ".weight_scale"); + VT_CHECK(ws.dtype == "F8_E4M3", + "qwen3_5 dense: expected F8_E4M3 weight_scale for " + proj); + const float ws2 = ReadF32Scalar(get(proj + ".weight_scale_2")); + VT_CHECK(ws2 != 0.0F, "qwen3_5 dense: zero weight_scale_2 for " + proj); + + Nvfp4Weight r; + r.n = out_dim; + r.k = in_dim; + r.scale2 = ws2; // ModelOpt stores the scale directly + r.weight_global_scale_inv = 1.0F / ws2; // the CT-convention divisor + // W4A16 unless the checkpoint also carries an activation scale. ModelOpt + // spells it `input_scale`; leaving alpha at 0 keeps IsTrueW4A4() false so the + // weight routes to the W4A16 dispatcher, matching vLLM's use_a16 branch. + // 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 float is = ReadF32Scalar(get(proj + ".input_scale")); + if (is != 0.0F) { + r.input_global_scale_inv = is; + r.alpha = r.scale2 * (1.0F / is); + } + } + r.packed = MakeOwned(vt::DType::kI8, {out_dim, in_dim / 2}); + VT_CHECK(packed.nbytes == r.packed.bytes.size(), + "qwen3_5 dense: ModelOpt packed byte-size mismatch for " + proj); + std::memcpy(r.packed.bytes.data(), packed.data, packed.nbytes); + MaybeReleaseSourcePages(packed.data, packed.nbytes); + r.scale = MakeOwned(vt::DType::kI8, {out_dim, in_dim / 16}); + VT_CHECK(ws.nbytes == r.scale.bytes.size(), + "qwen3_5 dense: ModelOpt scale byte-size mismatch for " + proj); + std::memcpy(r.scale.bytes.data(), ws.data, ws.nbytes); + MaybeReleaseSourcePages(ws.data, ws.nbytes); + return r; +} + GdnLayerWeights LoadGdnDense(const TensorResolver& get, const TensorExists& has, const std::string& base) { const std::string la = base + "linear_attn."; @@ -325,14 +393,31 @@ GdnLayerWeights LoadGdnDense(const TensorResolver& get, const TensorExists& has, // order (the checkpoint's in_proj_qkv already stacks q|k|v; z appended) and // in_proj_ba in exact [b,a] row order. The rollback paths take non-owning // row slices of these owners, so the split fields deliberately stay empty. - g.in_proj_qkvz = LoadMergedBf16RawNK( - get, {la + "in_proj_qkv.weight", la + "in_proj_z.weight"}); + // FP8 checkpoints (nvidia/Qwen3.6-27B-NVFP4 is `modelopt_mixed`) keep the GDN + // shards NATIVE. Merging forces one dtype, and dequantizing fp8 -> bf16 DOUBLES + // this tower's resident bytes (6.72 -> 13.44 GiB measured), which both slows + // decode and starves the KV pool. `ProjectGdnQkvz` already carries the + // separate-fp8 arm the 35B runs and selects it when the merged owner is empty. + const StTensor& qkv_probe = get(la + "in_proj_qkv.weight"); + if (qkv_probe.dtype == "F8_E4M3") { + g.in_proj_qkv_fp8 = LoadFp8RawShared(get, la + "in_proj_qkv"); + g.in_proj_z_fp8 = LoadFp8RawShared(get, la + "in_proj_z"); + } else { + g.in_proj_qkvz = LoadMergedBf16RawNK( + get, {la + "in_proj_qkv.weight", la + "in_proj_z.weight"}); + } g.in_proj_ba = LoadMergedBf16RawNK( get, {la + "in_proj_b.weight", la + "in_proj_a.weight"}); // NVFP4 checkpoints use compressed tensors; ordinary checkpoints use raw // torch-Linear BF16 [N,K]. - if (has(la + "out_proj.weight_packed")) { - g.out_proj_fp4 = LoadCtNvfp4Raw(get, la + "out_proj"); + if (IsNvfp4Projection(has, la + "out_proj")) { + g.out_proj_fp4 = LoadNvfp4AnyNaming(get, has, la + "out_proj"); + } else if (get(la + "out_proj.weight").dtype == "F8_E4M3") { + // Same rule as the in_proj shards above, and for the same measured reason: + // the bf16 arm dequantizes this tower and then runs it as a cuBLAS `gemvx`, + // which the decode profile shows costing far more than the bytes justify. + // `ProjectGdnOut` already selects `out_proj_fp8` when it is populated. + g.out_proj_fp8 = LoadFp8RawShared(get, la + "out_proj"); } else { g.out_proj = LoadBf16RawNK(get, la + "out_proj.weight"); } @@ -353,17 +438,27 @@ FullAttnLayerWeights LoadAttnDense(const TensorResolver& get, const std::string& base) { const std::string sa = base + "self_attn."; FullAttnLayerWeights a; + // Three forms, not two. `modelopt_mixed` checkpoints quantize this tower to + // FP8 W8A8 while leaving the MLP NVFP4, and a projection that matches neither + // the NVFP4 probe nor an FP8 dtype is genuinely BF16. Without the middle + // branch an FP8 tower fell through to `LoadBf16RawNK`, which dequantizes it: + // 1.562 GiB of FP8 became 3.12 GiB of BF16 re-read every decode step and + // executed as cuBLAS `gemvx`. The `*_fp8` slots and their `MatmulFp8Cutlass*` + // consumers already exist and are what the 35B runs. const auto load_projection = [&](const std::string& name, Nvfp4Weight& fp4, - OwnedTensor& plain) { - if (has(name + ".weight_packed")) - fp4 = LoadCtNvfp4Raw(get, name); - else + Fp8Weight& fp8, OwnedTensor& plain) { + if (IsNvfp4Projection(has, name)) { + fp4 = LoadNvfp4AnyNaming(get, has, name); + } else if (get(name + ".weight").dtype == "F8_E4M3") { + fp8 = LoadFp8RawShared(get, name); + } else { plain = LoadBf16RawNK(get, name + ".weight"); + } }; - load_projection(sa + "q_proj", a.q_proj_fp4, a.q_proj); - load_projection(sa + "k_proj", a.k_proj_fp4, a.k_proj); - load_projection(sa + "v_proj", a.v_proj_fp4, a.v_proj); - load_projection(sa + "o_proj", a.o_proj_fp4, a.o_proj); + load_projection(sa + "q_proj", a.q_proj_fp4, a.q_proj_fp8, a.q_proj); + load_projection(sa + "k_proj", a.k_proj_fp4, a.k_proj_fp8, a.k_proj); + load_projection(sa + "v_proj", a.v_proj_fp4, a.v_proj_fp8, a.v_proj); + load_projection(sa + "o_proj", a.o_proj_fp4, a.o_proj_fp8, a.o_proj); a.q_norm = LoadModelBf16Direct(get, sa + "q_norm.weight"); a.k_norm = LoadModelBf16Direct(get, sa + "k_norm.weight"); return a; @@ -374,10 +469,10 @@ DenseMlpWeights LoadDenseMlp(const TensorResolver& get, const TensorExists& has, const std::string& base) { const std::string mlp = base + "mlp."; DenseMlpWeights m; - if (has(mlp + "gate_proj.weight_packed")) { - m.gate_proj_fp4 = LoadCtNvfp4Raw(get, mlp + "gate_proj"); - m.up_proj_fp4 = LoadCtNvfp4Raw(get, mlp + "up_proj"); - m.down_proj_fp4 = LoadCtNvfp4Raw(get, mlp + "down_proj"); + if (IsNvfp4Projection(has, mlp + "gate_proj")) { + m.gate_proj_fp4 = LoadNvfp4AnyNaming(get, has, mlp + "gate_proj"); + m.up_proj_fp4 = LoadNvfp4AnyNaming(get, has, mlp + "up_proj"); + m.down_proj_fp4 = LoadNvfp4AnyNaming(get, has, mlp + "down_proj"); } else { m.gate_up_proj = dense_loaders::LoadMergedBf16RawNK( get, {mlp + "gate_proj.weight", mlp + "up_proj.weight"}); diff --git a/src/vllm/model_executor/models/qwen3_5_weights.cpp b/src/vllm/model_executor/models/qwen3_5_weights.cpp index 4276867e2..de73a1265 100644 --- a/src/vllm/model_executor/models/qwen3_5_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_weights.cpp @@ -445,6 +445,11 @@ Qwen3_5MoeLayerWeights LoadLayerImpl(const TensorResolver& get, } // namespace +// External-linkage seam so the DENSE loader can keep an FP8 GDN tower native. +Fp8Weight LoadFp8RawShared(const TensorResolver& get, const std::string& proj) { + return LoadFp8Raw(get, proj); +} + Qwen3_5MoeLayerWeights LoadQwen3_5MoeLayer(const TensorResolver& get, const std::string& layer_type, int64_t layer_idx, diff --git a/src/vt/cuda/cuda_moe.cu b/src/vt/cuda/cuda_moe.cu index 543ad793b..f24d539fa 100644 --- a/src/vt/cuda/cuda_moe.cu +++ b/src/vt/cuda/cuda_moe.cu @@ -146,26 +146,46 @@ __global__ void MoeRouterTopKKernel(float* weights, int32_t* indices, const Tin* li = static_cast(idx); } } - red[threadIdx.x] = lv; - redi[threadIdx.x] = li; - __syncthreads(); - for (int s = kBlock / 2; s > 0; s /= 2) { - if (static_cast(threadIdx.x) < s) { - const float ov = red[threadIdx.x + s]; - const int oi = redi[threadIdx.x + s]; - const float cv = red[threadIdx.x]; - const int ci = redi[threadIdx.x]; - // Higher value wins; on an exact tie the lower expert index wins. - if (ov > cv || (ov == cv && oi >= 0 && (ci < 0 || oi < ci))) { - red[threadIdx.x] = ov; - redi[threadIdx.x] = oi; - } + // Warp-shuffle argmax, then ONE cross-warp pass. This replaces a + // block-wide tree that cost log2(kBlock) __syncthreads PER ROUND: at + // decode the grid is one block per token, so a k=8 top-k over 256 experts + // spent ~64 barriers on a single SM and the measured kernel was 19.5 us + // for work that is a few hundred comparisons (4.7% of the 35B decode + // step). Barriers per round drop from log2(kBlock)+1 to 2. + // + // BYTE-IDENTICAL, and the reason is worth stating: this is an ARGMAX + // reduction, not an arithmetic one. The comparison below is exactly the + // one the tree used -- higher value wins, and on an exact tie the lower + // expert index wins -- and argmax under a total order is associative and + // commutative, so ANY reduction order yields the same (value, index). + // The softmax max and sum reductions above are arithmetic and their tree + // is deliberately left untouched, because changing THEIR order would + // change the denominator in the last ulp. + for (int off = 16; off > 0; off >>= 1) { + const float ov = __shfl_down_sync(0xffffffffu, lv, off); + const int oi = __shfl_down_sync(0xffffffffu, li, off); + if (ov > lv || (ov == lv && oi >= 0 && (li < 0 || oi < li))) { + lv = ov; + li = oi; } - __syncthreads(); } - const float best_v = red[0]; - const int best = redi[0]; + constexpr int kWarps = kBlock / 32; + if ((threadIdx.x & 31u) == 0u) { + red[threadIdx.x >> 5] = lv; + redi[threadIdx.x >> 5] = li; + } + __syncthreads(); if (threadIdx.x == 0) { + float best_v = red[0]; + int best = redi[0]; + for (int w = 1; w < kWarps; ++w) { + const float ov = red[w]; + const int oi = redi[w]; + if (ov > best_v || (ov == best_v && oi >= 0 && (best < 0 || oi < best))) { + best_v = ov; + best = oi; + } + } if (best >= 0) sp[best] = -INFINITY; // exclude from subsequent rounds weights[row * k + j] = best_v; indices[row * k + j] = static_cast(best); @@ -207,6 +227,7 @@ __global__ void MoeRouterGroupedTopKKernel(float* weights, int32_t* indices, float* gscore = smem + 2 * e; // [n_group] float* gkeep = smem + 2 * e + n_group; // [n_group] 0/1 mask __shared__ float red[kBlock]; + __shared__ int redi[kBlock]; // argmax partner for red, step (4) // (1) scores = softmax(logits, -1) | sigmoid(logits) (:110-117) if (sigmoid) { @@ -256,9 +277,11 @@ __global__ void MoeRouterGroupedTopKKernel(float* weights, int32_t* indices, } __syncthreads(); - if (threadIdx.x != 0) return; - + // Steps (2) and (3) stay serial on thread 0: they are O(n_group * group_size) + // ONCE per token and n_group is 1 or 8 for the shapes we serve. The block is + // NOT retired here any more, because step (4) below now uses it. const int64_t group_size = e / n_group; + if (threadIdx.x == 0) { // Group score: top-2 SUM with a bias (:124-126), else the group MAX (:128-131). for (int64_t g = 0; g < n_group; ++g) { const int64_t base = g * group_size; @@ -298,28 +321,72 @@ __global__ void MoeRouterGroupedTopKKernel(float* weights, int32_t* indices, if (best < 0) break; // fewer groups than topk_group (wrapper forbids it) gkeep[best] = 1.0f; } - for (int64_t g = 0; g < n_group; ++g) { - if (gkeep[g] != 0.0f) continue; - for (int64_t j = 0; j < group_size; ++j) sel[g * group_size + j] = -INFINITY; + for (int64_t g = 0; g < n_group; ++g) { + if (gkeep[g] != 0.0f) continue; + for (int64_t j = 0; j < group_size; ++j) sel[g * group_size + j] = -INFINITY; + } } + __syncthreads(); // the group mask in sel[] must be visible to the block + // (4) top-k over the masked selection scores; weight from the unbiased score. - float denom = 0.0f; + // + // BLOCK-PARALLEL, and byte-identical to the serial scan it replaces. This was + // `for idx in [0,e)` on THREAD 0 alone, k times: at Kimi-Linear's shape + // (e=256, k=8) that is ~2048 serial compares on one lane of one block, which + // the kernel's own comment deferred to "W9, after the numerics are gated". + // The identity argument is the same one used for the ungrouped router: this + // is an ARGMAX over a total order (higher score wins, exact tie to the lower + // expert index, which the serial ascending scan with strict `>` also gives), + // and argmax is associative and commutative, so any reduction order returns + // the same (value, index). Everything ARITHMETIC is untouched and still runs + // on thread 0 in the same order: the `denom` accumulation in k order, the + // renormalize, and the routed_scaling_factor. + float denom = 0.0f; // meaningful on thread 0 only + constexpr int kWarps = kBlock / 32; for (int j = 0; j < k; ++j) { - int64_t best = -1; - float best_v = -INFINITY; - for (int64_t idx = 0; idx < e; ++idx) { - if (sel[idx] > best_v) { - best_v = sel[idx]; - best = idx; + float lv = -INFINITY; + int li = -1; + for (int64_t idx = threadIdx.x; idx < e; idx += blockDim.x) { + const float v = sel[idx]; + if (v > lv) { // strict `>`, ascending stride -> lowest index at the max + lv = v; + li = static_cast(idx); + } + } + for (int off = 16; off > 0; off >>= 1) { + const float ov = __shfl_down_sync(0xffffffffu, lv, off); + const int oi = __shfl_down_sync(0xffffffffu, li, off); + if (ov > lv || (ov == lv && oi >= 0 && (li < 0 || oi < li))) { + lv = ov; + li = oi; + } + } + if ((threadIdx.x & 31u) == 0u) { + red[threadIdx.x >> 5] = lv; + redi[threadIdx.x >> 5] = li; + } + __syncthreads(); + if (threadIdx.x == 0) { + float best_v = red[0]; + int best = redi[0]; + for (int w2 = 1; w2 < kWarps; ++w2) { + const float ov = red[w2]; + const int oi = redi[w2]; + if (ov > best_v || (ov == best_v && oi >= 0 && (best < 0 || oi < best))) { + best_v = ov; + best = oi; + } } + if (best < 0) best = 0; // all -inf: the serial scan also fell back to 0 + sel[best] = -INFINITY; + const float w = orig[best]; + weights[row * k + j] = w; + indices[row * k + j] = static_cast(best); + denom += w; } - if (best < 0) best = 0; - sel[best] = -INFINITY; - const float w = orig[best]; - weights[row * k + j] = w; - indices[row * k + j] = static_cast(best); - denom += w; + __syncthreads(); // sel[best]=-INF visible + red/redi reusable next round } + if (threadIdx.x != 0) return; // (5) renormalize (:156-157) THEN routed_scaling_factor (:159-160). if (renormalize) { if (!(denom > 0.0f)) denom = 1.0f; diff --git a/tests/parity/hf_snapshot.h b/tests/parity/hf_snapshot.h new file mode 100644 index 000000000..336ce9695 --- /dev/null +++ b/tests/parity/hf_snapshot.h @@ -0,0 +1,66 @@ +#ifndef VLLM_TESTS_PARITY_HF_SNAPSHOT_H_ +#define VLLM_TESTS_PARITY_HF_SNAPSHOT_H_ + +// Resolving a Hugging Face cache snapshot for a checkpoint-gated test. +// +// Every one of these gates used to take the FIRST entry `directory_iterator` +// yielded under `/snapshots/`. That is only safe while a repo has exactly +// one cached revision, and `unsloth/Qwen3.6-27B-NVFP4` does not: +// +// @890bdef7 genuine NVFP4 - `weight_packed` U8 + `weight_scale` F8_E4M3 + +// `weight_global_scale` F32. Every committed 27B golden was +// captured against it (see the `oracle.model` field of +// tests/parity/goldens/qwen36_*_27b/manifest.json). +// @ccdaab7e the SAME repo name, silently re-quantized to FP8 W8A8 +// throughout, with every NVFP4-specific `*_global_scale` tensor +// gone. +// +// So the filesystem decided which model the SACRED gate measured, and a +// token-exact pass against an FP8 model would have been recorded as an NVFP4 +// pass. Publishers re-quantize in place; a correctness gate must name the +// revision its golden belongs to. + +#include +#include +#include +#include + +namespace parity { + +// The revision the committed 27B goldens were captured against. +inline constexpr const char* kQwen27NvfP4Revision = + "890bdef7a42feba6d83b6e17a03315c694112f2a"; + +// Snapshot directory for `` at `revision`, or "" when it is not cached +// (the caller then emits its loud SKIP). `env_override`, when set and non-empty, +// names an explicit snapshot directory for a deliberate different-checkpoint +// run and is the ONLY way to gate a revision other than the pinned one -- a +// cache holding some other revision skips rather than silently substituting it. +inline std::string HfSnapshot(const char* repo_dir, const char* revision, + const char* env_override) { + namespace fs = std::filesystem; + std::error_code ec; + if (env_override != nullptr) { + const char* over = std::getenv(env_override); + if (over != nullptr && *over != '\0') { + if (fs::exists(fs::path(over) / "config.json", ec)) return over; + return ""; + } + } + const char* home = std::getenv("HOME"); + if (home == nullptr) return ""; + const fs::path snap = fs::path(home) / ".cache/huggingface/hub" / repo_dir / + "snapshots" / revision; + if (!fs::exists(snap / "config.json", ec)) return ""; + return snap.string(); +} + +// The 27B NVFP4 gate model, pinned to the goldens' revision. +inline std::string Qwen27NvfP4Snapshot() { + return HfSnapshot("models--unsloth--Qwen3.6-27B-NVFP4", + kQwen27NvfP4Revision, "VT_QWEN27_SNAPSHOT"); +} + +} // namespace parity + +#endif // VLLM_TESTS_PARITY_HF_SNAPSHOT_H_ diff --git a/tests/parity/test_op_parity.cpp b/tests/parity/test_op_parity.cpp index 3ba721292..a2b3f303c 100644 --- a/tests/parity/test_op_parity.cpp +++ b/tests/parity/test_op_parity.cpp @@ -17,6 +17,7 @@ #include #include "npy.h" +#include "hf_snapshot.h" #include "vllm/model_executor/layers/rotary_embedding/base.h" #include "vllm/model_executor/model_loader/safetensors_reader.h" #include "vllm/model_executor/models/model_registry.h" @@ -1217,17 +1218,8 @@ bool RunQwen36Logits(Backend& /*b*/, Queue& q, const fs::path& dir, // Snapshot dir of the 27B dense checkpoint (contains config.json), or "". std::string Find27BSnapshot() { - const char* home = std::getenv("HOME"); - if (home == nullptr) return ""; - const fs::path snaps = - fs::path(home) / - ".cache/huggingface/hub/models--unsloth--Qwen3.6-27B-NVFP4/snapshots"; - std::error_code ec; - if (!fs::is_directory(snaps, ec)) return ""; - for (const auto& e : fs::directory_iterator(snaps, ec)) { - if (fs::exists(e.path() / "config.json", ec)) return e.path().string(); - } - return ""; + // Pinned to the goldens' revision; see tests/parity/hf_snapshot.h. + return parity::Qwen27NvfP4Snapshot(); } // --- 27B DENSE full-model logits acceptance gate. diff --git a/tests/parity/test_qwen27_ngram_spec_decode.cpp b/tests/parity/test_qwen27_ngram_spec_decode.cpp index 4583ddc7c..494d22da1 100644 --- a/tests/parity/test_qwen27_ngram_spec_decode.cpp +++ b/tests/parity/test_qwen27_ngram_spec_decode.cpp @@ -35,6 +35,8 @@ #include +#include "hf_snapshot.h" + #include "vllm/config/speculative.h" #include "vllm/entrypoints/model_loader.h" #include "vllm/sampling_params.h" @@ -48,20 +50,8 @@ namespace { // matching the ngram golden + the 27B SACRED gate), preferring the single // model.safetensors snapshot over a sharded FP8-lm_head re-quant. std::string Snap27B() { - const char* home = std::getenv("HOME"); - if (home == nullptr) return ""; - const fs::path base = - fs::path(home) / - ".cache/huggingface/hub/models--unsloth--Qwen3.6-27B-NVFP4/snapshots"; - std::error_code ec; - if (!fs::is_directory(base, ec)) return ""; - std::string any; - for (const auto& e : fs::directory_iterator(base, ec)) { - if (!fs::exists(e.path() / "config.json", ec)) continue; - any = e.path().string(); - if (fs::exists(e.path() / "model.safetensors", ec)) return e.path().string(); - } - return any; + // Pinned to the goldens' revision; see tests/parity/hf_snapshot.h. + return parity::Qwen27NvfP4Snapshot(); } vllm::SamplingParams Greedy(int max_tokens) { diff --git a/tests/parity/test_qwen27_paged_engine.cpp b/tests/parity/test_qwen27_paged_engine.cpp index b3ae65f30..ec701fffd 100644 --- a/tests/parity/test_qwen27_paged_engine.cpp +++ b/tests/parity/test_qwen27_paged_engine.cpp @@ -34,6 +34,7 @@ #endif #include "npy.h" +#include "hf_snapshot.h" #include "vllm/entrypoints/model_loader.h" #include "vllm/model_executor/models/qwen3_5_internal.h" #include "vllm/sampling_params.h" @@ -74,21 +75,12 @@ void CheckDeviceCacheResidency( #endif } -// Snapshot dir of the 27B checkpoint (contains config.json), or "". Same HF -// cache layout as Find35BSnapshot, for models--unsloth--Qwen3.6-27B-NVFP4. -std::string Find27BSnapshot() { - const char* home = std::getenv("HOME"); - if (home == nullptr) return ""; - const fs::path snaps = - fs::path(home) / - ".cache/huggingface/hub/models--unsloth--Qwen3.6-27B-NVFP4/snapshots"; - std::error_code ec; - if (!fs::is_directory(snaps, ec)) return ""; - for (const auto& e : fs::directory_iterator(snaps, ec)) { - if (fs::exists(e.path() / "config.json", ec)) return e.path().string(); - } - return ""; -} +// Snapshot dir of the 27B gate checkpoint, or "" to SKIP. Pinned to the +// revision the committed goldens were captured against -- this repo has TWO +// cached snapshots under one name and only one of them is NVFP4. See +// tests/parity/hf_snapshot.h for why an unpinned lookup could gate the wrong +// weights. +std::string Find27BSnapshot() { return parity::Qwen27NvfP4Snapshot(); } // Load an i32 (.npy " LoadI32Npy(const fs::path& p) { diff --git a/tests/parity/test_qwen27_spec_decode.cpp b/tests/parity/test_qwen27_spec_decode.cpp index 129822ee7..4c7dd5741 100644 --- a/tests/parity/test_qwen27_spec_decode.cpp +++ b/tests/parity/test_qwen27_spec_decode.cpp @@ -34,6 +34,7 @@ #include #include "npy.h" +#include "hf_snapshot.h" #include "vllm/config/speculative.h" #include "vllm/entrypoints/model_loader.h" #include "vllm/sampling_params.h" @@ -43,17 +44,8 @@ namespace fs = std::filesystem; namespace { std::string Find27BSnapshot() { - const char* home = std::getenv("HOME"); - if (home == nullptr) return ""; - const fs::path snaps = - fs::path(home) / - ".cache/huggingface/hub/models--unsloth--Qwen3.6-27B-NVFP4/snapshots"; - std::error_code ec; - if (!fs::is_directory(snaps, ec)) return ""; - for (const auto& e : fs::directory_iterator(snaps, ec)) { - if (fs::exists(e.path() / "config.json", ec)) return e.path().string(); - } - return ""; + // Pinned to the goldens' revision; see tests/parity/hf_snapshot.h. + return parity::Qwen27NvfP4Snapshot(); } std::vector LoadI32Npy(const fs::path& p) { diff --git a/tests/parity/test_qwen27_spec_decode_concurrent.cpp b/tests/parity/test_qwen27_spec_decode_concurrent.cpp index 877785038..0b9928bdb 100644 --- a/tests/parity/test_qwen27_spec_decode_concurrent.cpp +++ b/tests/parity/test_qwen27_spec_decode_concurrent.cpp @@ -30,6 +30,8 @@ #include #include +#include "hf_snapshot.h" + #include "vllm/config/speculative.h" #include "vllm/entrypoints/model_loader.h" #include "vllm/model_executor/models/qwen3_5.h" @@ -40,17 +42,8 @@ namespace fs = std::filesystem; namespace { std::string Find27BSnapshot() { - const char* home = std::getenv("HOME"); - if (home == nullptr) return ""; - const fs::path snaps = - fs::path(home) / - ".cache/huggingface/hub/models--unsloth--Qwen3.6-27B-NVFP4/snapshots"; - std::error_code ec; - if (!fs::is_directory(snaps, ec)) return ""; - for (const auto& e : fs::directory_iterator(snaps, ec)) { - if (fs::exists(e.path() / "config.json", ec)) return e.path().string(); - } - return ""; + // Pinned to the goldens' revision; see tests/parity/hf_snapshot.h. + return parity::Qwen27NvfP4Snapshot(); } vllm::SamplingParams Greedy(int max_tokens) {