diff --git a/.agents/issue-index.md b/.agents/issue-index.md index ff108e7fd..10d34c047 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -733,6 +733,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1933](https://github.com/mudler/vllm.cpp/issues/1933) | `MODEL-TEXT-deepseek-v4-deepseek-v4-for-causal-lm` | **The GGUF arm resolved `deepseek-llm`, `deepseek-v3` and `joyai-llm` to `kLlama3` as a documented "close APPROXIMATION" while the exact pre-tokenizer for each was in the tree.** Found closing [#1924](https://github.com/mudler/vllm.cpp/issues/1924), and it contradicts that issue's own scope note ("The GGUF path is unaffected because it carries its own vocabulary") — the vocabulary is its own, the pre-tokenizer was not. `deepseek-llm` is `LLAMA_VOCAB_PRE_TYPE_DEEPSEEK_LLM` = `kDeepSeek`, which landed at `66a44f9bf` and was never wired to the pre name; `deepseek-v3` and `joyai-llm` (plus `hunyuan-dense`, which was refused by name entirely) are `LLAMA_VOCAB_PRE_TYPE_DEEPSEEK3_LLM` = the `kDeepSeekV3` #1924 adds. NOT a rare-boundary difference: the V3 alternation binds an ASCII punctuation character to the letters after it, so `def foo(x): return x` keeps `(x` as one piece where `kLlama3` splits it, and `$var`/`_name` are one piece against two; `kDeepSeek` isolates every newline and splits digits one at a time against `kLlama3`'s groups of three, so every multi-digit number in a prompt got a different id. Same shape as [#347](https://github.com/mudler/vllm.cpp/issues/347). The artifact it bites is the DeepSeek-V4-Flash GGUF (`antirez/ds4` q2-imatrix, pre `joyai-llm`), which the old comment named by hand. `laguna` KEEPS `kLlama3` and its approximation note: llama.cpp has no `laguna` pre name, so nothing exact exists to resolve it onto, and that distinction is pinned in the test rather than left to the reader. FIXED IN FLOW with #1924, because the exact V3 pipeline half of it needs lands in the same change | bug | | [#1904](https://github.com/mudler/vllm.cpp/issues/1904) | `LTX25-VAE-DEVICE-RESIDENCY` | **The LTX-2.5 video VAE hand-rolls `DevBuf` instead of the shared `dense_attn::DBuf` device-buffer seam.** W5 ([#1007](https://github.com/mudler/vllm.cpp/issues/1007)) added `DevBuf` at `src/vllm/model_executor/models/ltx2_video_vae.cpp:145-170` — move-deleted RAII over `vt::Backend::Alloc`/`Copy`/`Free` with a `Download` helper — which is a second copy of `vllm::dense_attn::DBuf` (`include/vllm/model_executor/models/dense_device_glue.h:109`), the same object with the same constructor shape and the same `.t()`/`.Download()` surface. `AGENTS.md` `## Shared seams` forbids a hand-written parallel path and no exception is recorded. The difference is not cosmetic: `DBuf` draws from the shared `DevicePool` (`device_pool.h:71`) so a block is reused, while `DevBuf` calls `Alloc`/`Free` directly and `Conv3dThroughSeam` builds three to four of them PER CONVOLUTION, so a decode performs a driver `Free` per operand per convolution on the one path the pool exists to serve. NOT fixed in the flow that found it, and the reason is a behaviour change rather than time: `DBuf` resolves `platforms::GetPlatform(device.type)` through `ResolveDevicePoolPolicy` (`dense_device_glue.h:88-105`) and THROWS for a device type whose platform was never registered, so the switch makes a registered platform a new precondition of a decode that has none today. The audit of whether any current caller reaches the video VAE on such a device is part of the issue. Found by `LTX25-VAE-DEVICE-RESIDENCY` while porting the decode onto a resident volume ([#1451](https://github.com/mudler/vllm.cpp/issues/1451)); listed under `## Owed` in [`ltx25-vae-device-residency.md`](specs/ltx25-vae-device-residency.md) | enhancement | | [#1939](https://github.com/mudler/vllm.cpp/issues/1939) | `LTX25-DEVICE-RESIDENCY` | **The `vt::Conv3d` and `vt::Conv1d` CUDA byte-identity cases score as PASSES on every CI lane, because a doctest `[SKIP]` is a pass and no lane here has a GPU.** Each begins `if (!HasCuda()) { printf("[SKIP] ..."); return; }`. MEASURED by the fresh review of [#1938](https://github.com/mudler/vllm.cpp/pull/1938): under `ctest --output-on-failure`, which is how `.github/workflows/ci.yml` invokes these binaries, stdout from a PASSING test is discarded, so the `[SKIP]` lines are never printed at all, and doctest reports `passed` with `0 skipped` because an early `return` from a case body is indistinguishable from a case with no failing assertion. Negative with its search: nothing under `scripts/` or `.github/` greps for `[SKIP]`, so no checker can tell an executed device case from a skipped one. The [#1452](https://github.com/mudler/vllm.cpp/issues/1452) measurement therefore exists only as a one-off taken by hand in an `rc` lease, and `docs/models/ltx-2-5.md` now says so rather than reading as continuous gating. This is the shape `src/vt/cuda/cuda_backend.cu:341-346` already names in prose, and a standing property of the suite rather than debt #1452 introduced, which is why it is filed rather than fixed there. TWO RIDERS from the same review. (1) The conv3d cancellation case's teeth check is COUNT-BASED and guards half of what it claims: `CHECK(differing > 0)` bites on the mutation its comment names — deleting the shared weight row drives `differing` to 0 and the check red — but NOT on losing the magnitude, where `kBig = 1.0` instead of 2^40 leaves `differing` at 2266 while the maximum absolute difference collapses from **9275.17** to **9.8e-4**, so the case loses its discriminating power and the guard still passes; a magnitude assertion with those two measured values as the separation closes both halves. (2) `RunCaseF32` leaks its device buffers and queue if `vt::Conv3d` throws, identical to the pre-existing `Stage()` in `tests/vt/test_ops_conv1d_general.cpp`, so a copied pattern rather than a regression, and it fires only on the refusal paths. The byte-identity loop also aborts on the first device exception, reporting one shape instead of the full picture. What would close it: a CI lane that owns a GPU, or an in-tree assertion distinguishing "the device arm ran" from "there was no device" — the tree has the idiom already, `vt::GetOpProviderStats` / `vt::OpProviderNameAt` as `tests/vt/test_ops_mamba2_ssd.cpp` uses it — red-first against a build with the CUDA arm deregistered. NOT fixed in flow: it needs its own red-before evidence and touches a second suite this row does not own. Listed under `## Owed` in [`ltx25-device-residency.md`](specs/ltx25-device-residency.md) | bug | +| [#1961](https://github.com/mudler/vllm.cpp/issues/1961) | `MODEL-DSV4-EXL3` | **DeepSeek-V4's doubled DSA tensors are the `coff=2` overlapping-window pair, and our host forward has no composition to put them in — plus `dsa_dense` rests on an exactness claim that is false at every sequence length.** SCOPED 2026-08-26 against the PRIMARY oracle vLLM at the parity pin `5559679229bc961848b121ccdeaa8fa5d79bec98` by [dsv4-dsa-geometry.md](specs/dsv4-dsa-geometry.md); no secondary oracle is used or needed, because vLLM registers and implements this architecture in full. **What the width is:** `compress_ratios` in the real artifact's `config.json` is a PER-LAYER list — `[0, 0, 4, 128, 4, 128, ..., 4, 0, 0, 0]`, giving 21 layers at `cr == 4`, 20 at `cr == 128` and 2 dense, which is the row's "41 of 43 carry a compressor, 21 carry an indexer" with its reason attached. Upstream turns that value into the width in one line (`vllm/models/deepseek_v4/compressor.py:247-248`: `self.overlap = compress_ratio == 4; self.coff = 1 + self.overlap`), spent at `:279-287` on `[coff*head_dim, coff*head_dim]` and at `:270-277` on an `ape` of `[compress_ratio, coff*head_dim]` — every measured width with no residue, and the `cr == 128` layers collapse to `coff == 1`, which is why 20 of the 41 already load. The two halves are **the two overlapping compression windows a token belongs to**: the pooling window is `coff*compress_ratio` wide while a row is emitted every `compress_ratio` tokens (`compressor.py:171-173`), so at `cr == 4` an 8-token window steps by 4 and every token is pooled twice, once in each role. The half is selected at GATHER time by window position and the weight is never split (`common/ops/fused_compress_quant_cache.py:182`, inside the main compressor's `_fused_kv_compress_norm_rope_insert_sparse_attn`: `head_offset = (tokens >= COMPRESS_RATIO) * HEAD_SIZE`, where `HEAD_SIZE` is `head_dim`, not the stored width — the indexer and mxfp4-indexer kernels carry the same line at `:730` and `:909`, so the construct does not pick out its own line); the only split `packed_modules_mapping` performs is the OPPOSITE one, merging the stored `wkv` and `wgate` into one GEMM (`nvidia/model.py:1157-1158`). Not a gate/value pair, not an interleave, not a fusion. **FOUR tensors refuse, not the three `MODEL-DSV4-EXL3` `## Owed` names:** `attn.compressor.ape` `[4, 1024]` vs our `{4, 512}` is missing from it, and `attn.indexer.wq_b.weight` `[8192, 1024]` is listed there as a width problem when it is a WRONG-INPUT-SPACE problem — upstream is `ReplicatedLinear(q_lora_rank, head_dim*n_head)` called on `qr` in `DeepseekV4Indexer.forward` (`attention.py:721-726`, `:835`), so `[8192, 1024]` is `[inh*ihd, q_lora_rank]` at natural size while we ask for `[inh*ihd, H]` and feed it `x`; no gate ever saw it because the collapsed fixture WRITES `wq_b` at `K = H` to match what our forward feeds it, so the two agree by construction — NOT because `H` and `q_lora_rank` coincide, which they do not (`dsv4_exl3_fixture.h:141,149`: `kHidden` 256, `kQLora` 128). #1970 repairs the loader half and adds `FixtureOptions::collapsed_indexer_wq_b`, the case that reaches the check; the forward still feeds `x`, so the input-space defect stands. **The composition gap:** our `AttentionBlock` (`deepseek_v4.cpp:827-857` at this branch's head; `:721-751` when this row was written) pools a fixed `win = 2` window of the MLA's own `kraw`, for EVERY token, overwriting the dense latent in place; upstream pools `coff*cr` rows of a SEPARATE `compressor.wkv` only at `(position+1) % cr == 0`, into a SEPARATE compressed KV cache beside a SWA(128) raw cache, with the indexer selecting among COMPRESSED rows. There is no half of these tensors our forward wants, because the composition they belong to is not there. **THE FINDING that reaches past this row:** `deepseek_v4.cpp:763-775` (`:664-676` when this row was written) justifies forcing DSA off with "dense MLA is EXACT ... whenever `seq_len <= index_topk` (=512)", and [#1925](https://github.com/mudler/vllm.cpp/issues/1925) quotes it onward — it is right about the indexer and WRONG about the attention. On a `cr > 1` layer ONE kernel takes ONE softmax over the UNION of the raw sliding window and the selected compressed rows (`nvidia/flashinfer_sparse.py:769-782`, in `DeepseekV4FlashInferSM120Attention._forward_decode`; the same call is at `:486`, `:511` and `:888`); compressed rows are POOLED AGGREGATES of `coff*cr` raw rows, so no selection over them reproduces attention over raw rows, and their count does not depend on `index_topk` — a 10-token prefill at `cr == 4` already has two, and the short-context branch explicitly still builds the K cache and selects all candidates (`attention.py:813-830`). Upstream at `seq_len == 10` attends 10 raw AND 2 compressed keys; we attend 10. Upstream's attention here is HIERARCHICAL — recent tokens at full resolution, older tokens pooled `cr:1`, jointly normalized — and dense causal attention is not that at ANY sequence length, so a token gate cannot detect it above or below 512. The GGUF arm runs `dsa_dense` on the real geometry today, so the shipping GGUF DeepSeek-V4 path is already not upstream's attention on 41 of 43 layers. Two riders: `dsa_dense = (be.gguf != nullptr)` keys off the WEIGHT SOURCE while upstream keys off `compress_ratios[layer_id]` (`attention.py:209`, `:274`, `:334`, `flashinfer_sparse.py:263` in `DeepseekV4FlashInferMLAAttention.forward_mqa`, repeated at `:686` and `:793`), and upstream's "dense" layers are SLIDING-WINDOW 128, not dense (`attention.py:204`), where our forward has no sliding window at all. **NOT FIXED IN FLOW, deliberately:** the loader half is small (derive widths as `coff = 1 + (cr == 4)`, take `wq_b`'s K from `q_lora_rank`) but landing only it is WORSE than the refusal it removes, because materializing `comp_wgate` at `[1024, 4096]` for a call with `hd == 512` mis-indexes it. WITHDRAWN AS WRITTEN, and #1970's row carries the same withdrawal: this said `Gemm`'s host arm is a `MatVec` with "no length check" and that the result is a silently wrong number. `deepseek_v4.cpp:413` is an unconditional `VT_CHECK` and `Gemm`'s keep-quant arm checks the shape too, so what the widened load without a refusal produces is an ANONYMOUS `vt: MatVec weight size mismatch`, not a wrong token. The refusal buys a DIAGNOSTIC and that is the whole of it. Three candidate shapes are set out in the spec (port upstream's DSA; a per-layer dense selector both arms read; loader-accepts/forward-refuses-by-name); choosing among them, and deciding what the row's equivalence gate compares against once the two arms stop sharing an attention path, is a design decision no helper owns and is returned as `NEEDS_DECISION`. Related cache topology: [#1925](https://github.com/mudler/vllm.cpp/issues/1925), [#1960](https://github.com/mudler/vllm.cpp/issues/1960). Listed under `## Owed` in [dsv4-dsa-geometry.md](specs/dsv4-dsa-geometry.md) | bug | | [#1946](https://github.com/mudler/vllm.cpp/issues/1946) | `SPEC-DFLASH2` | **The DFlash2 draft uploaded a SECOND device copy of the target's embedding table — BF16 `[248320, 5120]` = 2,542,796,800 B (2.543 GB) — because `ResidentWeight` caches its upload on the `OwnedTensor` and the draft held its own.** W9 ([#1849](https://github.com/mudler/vllm.cpp/issues/1849)) made both HOST reads borrow-first and scoped itself to the host in its own comment at `src/vllm/entrypoints/model_loader.cpp:358-360`; the `if (!w.d_dev)` guard at `include/vllm/model_executor/models/dense_attn_block.h:191` is per-tensor, so two `OwnedTensor`s meant two `d_dev` allocations of identical bytes whatever the host residency was. Upstream rebinds the MODULE by reference instead (`vllm/v1/worker/gpu/spec_decode/dflash/utils.py:64-74 @ b389ac29465b33f9e9c534df221ea3c129e9793f`, `del draft_inner.embed_tokens; draft_inner.embed_tokens = target_embed`) and holds one, which our own MTP lane already mirrors (`Qwen3_5MTPModel` points at the target's tensor) and the DFlash lane did not. GB10 is unified memory, so the second copy is 2.543 GB of the same 119 GiB the KV pool comes out of. Fixed in flow: the draft and the target now share ONE `OwnedTensor`, rebound at the one `LoadedEngine` constructor all three draft loaders cross. The `lm_head` half stays owed to the parent spec's `## Owed` O3. See [the embed device dedup spec](specs/dflash2-embed-device-dedup.md) | bug | | [#1951](https://github.com/mudler/vllm.cpp/issues/1951) | — | **The DSpark draft takes the SAME second device copy of the target's embedding table that [#1946](https://github.com/mudler/vllm.cpp/issues/1946) removed from the DFlash lane, whenever its checkpoint omits one.** `LoadDsparkDraft` moves the target's table into `draft->dspark->backbone.embed_tokens`, which is a second `OwnedTensor`, and `ResidentWeight` caches its device upload on the `OwnedTensor` itself (`include/vllm/model_executor/models/dense_attn_block.h::ResidentWeight`) — so it is a second device allocation of identical bytes, the exact defect #1946 measured at 2,542,796,800 B on the 27B. NOT fixed in flow, and the reason is structural rather than scheduling: `BindDflashDraftSharedEmbed` works because `Qwen3DFlashWeights` can carry a BORROWED `const OwnedTensor*` beside its own table, while the DSpark backbone owns its table BY VALUE inside `Qwen3DSparkWeights`, so rebinding `draft.weights.embed_tokens` there would touch a field the DSpark forward never reads and leave the copy that costs the memory in place. The skip is by name (`if (draft.dspark != nullptr) return false;`) and `tests/vllm/v1/spec_decode/test_dflash2_embed_dedup.cpp` pins it, so the gap cannot become silent. Both published DSpark drafts SHIP their own table, so nothing on the default published path duplicates today. Owed under `## Owed` O2 of [the embed device dedup spec](specs/dflash2-embed-device-dedup.md) | bug | | [#1953](https://github.com/mudler/vllm.cpp/issues/1953) | `SPEC-DFLASH2` | **`dense_attn::ResidentWeight` had no guard against an EMPTY tensor, so a cleared weight reached a kernel as a null host alias or as a zero-byte device allocation viewed at full shape.** Found by the fresh review of [#1952](https://github.com/mudler/vllm.cpp/pull/1952) while checking a claim that turned out to be false: three places justified [#1946](https://github.com/mudler/vllm.cpp/issues/1946)'s clear of the draft's own `embed_tokens` by saying a later read of that field would get an empty table "which `vt::Embedding` refuses by name rather than silently re-uploading 2.5 GB". It does not. `vt::Embedding` (`src/vt/ops.cpp`) validates ranks, shapes, dtypes, contiguity and device and NEVER the data pointer or the byte length, and `ResidentWeight` takes the shape from the CALLER, so an emptied tensor satisfies every `VT_CHECK` on the way down. The outcome is worse than the duplicate upload the clear prevents: the `is_cpu()` arm aliases a null host pointer into a kernel (SIGSEGV) and a device arm reads `bytes.size()` as 0, calls `d.b.Alloc(0)` and returns a `[vocab, H]` view over a zero-byte allocation — out-of-bounds device reads, which IS the silently-wrong-tokens failure the clear exists to stop. Measured under mutation, `REQUIRE(t.data != nullptr)` passed too, because a zero-size `Alloc` returns a valid one-byte pointer. NOT DFlash-specific: `ResidentWeight` is the shared device-residency seam, 373 call sites across 34 model files plus `include/vllm/model_executor/layers/linear.h`, and the tree's convention of guarding with `!Empty()` at the call site (`opt.cpp`'s `affine`, `phi.cpp`'s `BiasedProj`, `deepseek_v2.cpp`'s router bias, `qwen3_5.cpp:8532`) relies on every caller remembering, with nothing enforcing it. FIXED IN FLOW with #1946, because the false claim and the missing check are one defect and one gate covers both: the seam now refuses an empty weight on both arms by name. The predicate is `bytes.empty()` and not `OwnedTensor::Empty()`, since a weight whose host buffer was reclaimed after upload (`host_released`) is populated and served by the `d_dev` branch; and the staging assert sits inside `if (!w.d_dev)`, so an already-resident weight re-read on the decode path pays nothing. Red-first by the last case of `tests/vllm/v1/spec_decode/test_dflash2_embed_dedup.cpp`. `ResidentWeightF32` has the same shape and stays owed | bug | @@ -740,7 +741,9 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1910](https://github.com/mudler/vllm.cpp/issues/1910) | `BACKEND-ROCM` | **`KQuantGemmK` strides 32 lanes over `nsb = K/256` superblocks, so half of every warp idles on three quarters of decode calls.** `rocm_grouped_gemm.hip:449` gives one warp each `(i,j)` output and runs `for (sb = lane; sb < nsb; sb += 32)` followed by a fixed 5-round `__shfl_down_sync` reduction. On a 4096-wide model `nsb` is 16, so lanes 16..31 execute nothing and the reduction runs anyway. Measured by instrumenting the launcher on `4b1154bc5`, `Ornith-1.5-9B-Q4_K_M` (dense `qwen35`), RX 9060 XT (gfx1200), ROCm 7.2.3: **195 of 259 decode (`m=1`) dispatches carry `nsb` = 16**, led by 128 calls at `n=12288 k=4096` and including the lm_head at `n=248320 k=4096` Q6_K, one per token and 20x wider than any other output. **Not every shape is affected and a fix must not regress those:** the two `k=12288` entries have `nsb` = 48 and pack all 32 lanes, so the defect tracks `K` rather than the kernel. Cost, profiled with `rocprofv3 --kernel-trace --stats` by differencing `--max-tokens 4` against `--max-tokens 36` over 32 tokens WITH [#1876](https://github.com/mudler/vllm.cpp/issues/1876) applied: `KQuantGemmK`'s three instantiations total **22.937 ms/token, 54.3% of decode GPU time** over 129 calls, and with `wvSplitKSml` the matmul family is 33.65 ms against llama.cpp `b10451` HIP's 21.229 ms `mul_mat_vec_q` on the identical workload — **12.4 ms of the 18.76 ms/token gap that remains once #1876 lands, about 66% of it**. Same defect class as #1876, one kernel downstream: a decomposition written for prefill shapes that starves at `m` = 1. NOT established and stated rather than implied: the 593 us/call instantiation is **not** attributed to a call site (lm_head is the obvious candidate but dispatches were never correlated against the profiler's per-kernel rows, and that correlation decides whether a fix targets lm_head or the general `nsb`=16 path); no fix is proposed and no speed claim is made; the profile was taken on the UNMERGED `row/ROCM-Q8K-QUANT-DECOMP` branch, so on `main` this kernel's share is smaller while its absolute cost is identical; the host was not idle at loadavg 2.2-2.3 although free VRAM was asserted above 13 GiB with no resident model process; one model, one prompt, batch 1, gfx1200 only; and the MoE path at `:547` uses the same kernel and was not measured | perf | | [#1870](https://github.com/mudler/vllm.cpp/issues/1870) | `BACKEND-ROCM` | **`VT_GGUF_KEEP_QUANT=0` is documented as a same-binary opt-out and is unreachable on a 16 GiB discrete ROCm card, failing with a raw allocator throw rather than a refusal that names the cause.** `engine-fatal: EngineCore busy loop threw: vt rocm: hipMalloc: out of memory` on **both** `Qwen3.6-14B-A3B-VibeForged-v2-Q4_K_M` (7.87 GiB) and `Ornith-1.5-9B-Q4_K_M` (5.23 GiB) on `4b1154bc5`, RX 9060 XT (gfx1200, 15.92 GiB), ROCm 7.2.3, `--device auto`, with free VRAM asserted above 13 GiB and no resident model process; both run normally on the keep-quant default at 13.0-13.1 and 18.4-18.6 tok/s. The OOM itself is arithmetic and expected — a Q4_K_M expands roughly 4x to bf16, so 5.23 GiB becomes about 20 GiB and does not fit — and three things around it are the defect. (1) `docs/ENVIRONMENT.md:94` reads "`0` disables it and expands to BF16" and states **no memory precondition**, so the documented behavior is unreachable on this class of board and the document is wrong by omission. (2) `AGENTS.md` requires an unreachable arm to refuse with a message naming the missing part, and `hipMalloc: out of memory` names neither the knob, nor the expansion, nor the budget required. (3) It **removes the same-binary A/B lever** that `AGENTS.md` requires before a performance result is accepted, which is not hypothetical: it blocked the keep-quant attribution [#1863](https://github.com/mudler/vllm.cpp/issues/1863) wanted, and it is why [#1876](https://github.com/mudler/vllm.cpp/issues/1876) had to carry its own `VT_ROCM_Q8K_BLOCK` lever instead. Related gap in the same area: `kMoeGroupedGemmBf16` is unregistered on ROCm (`rocm_ops.hip` has zero occurrences, CUDA has it), so even where memory allowed the expansion the bf16 MoE arm has no provider. Split out of [#1506](https://github.com/mudler/vllm.cpp/issues/1506), whose title claim stopped being true when [#523](https://github.com/mudler/vllm.cpp/pull/523) registered `kMatmulBTQuant` on ROCm on 2026-08-21; its surviving `1.73x peak RSS` finding is this, re-measured, and on this card the penalty is no longer a ratio but a refusal to run. Filed separately rather than by re-scoping that issue, because the index is append-only and an edited row is duplicated rather than merged. A fix shape is bounded but NOT designed here: resolve the expanded residency requirement at load, compare against the device budget, and refuse by name before allocating; whether the knob should instead be ignored with a warning is a product decision this row does not settle | bug | | [#1914](https://github.com/mudler/vllm.cpp/issues/1914) | `ENG-WEIGHT-OFFLOAD` | **Four measured ROCm device facts for the weight-offload row, from a throwaway gfx1200 spike that was never merged.** The row mirrors vLLM's `cpu_offload_gb` ([#797](https://github.com/mudler/vllm.cpp/issues/797), the dense half of [#149](https://github.com/mudler/vllm.cpp/issues/149)) and its config surface has landed, but **none of it has been measured on AMD** and `specs/weight-offload-uva.md`'s scope table names no ROCm arm. Spike `5056bbf90` on `spike/rocm-523`, 2026-08-19, base `7b9e207b1`, RX 9060 XT (gfx1200, 15.92 GiB), ROCm 7.2.3, +86 lines across three files, inert with no environment variable set so the OFF arm is the unmodified upload path in the same binary. (1) **The premise works:** `Qwen3.6-35B-A3B-UD-Q4_K_S` at 19.45 GiB dies on `hipMalloc: out of memory` and, with large weights kept host-resident and handed to the kernel as a device-readable pointer, loads and generates; on the 14B, offloading 2.00 of 6.39 GiB of experts (31%) gave **byte-identical tokens** for **10.6%**. No new backend virtual was needed for the pinned arm, because on ROCm `hipHostMalloc` returns a pointer the device reads directly and `hipHostGetDevicePointer` returns the SAME value. (2) **A slab-read microbenchmark overpredicts by about 3x:** 23 GB/s idealised streaming against roughly 8 for the real GEMM, and at 47% offloaded the same 60 GEMM dispatches went 8.03 -> 31.20 ms/token, so an offload budget sized from a streaming-bandwidth number will be optimistic. (3) **The budget is NOT monotonic:** 6 GiB gives 3.21 tok/s and 7 GiB gives 7.67, because 6 leaves almost nothing for KV and allocator slack — a SMALLER budget is 2.4x slower, the cliff is reproducible, and its mechanism is **unexplained**, so a naive "offload as little as possible" policy walks into it. (4) **`hipMallocManaged` does NOT migrate on this part:** it allocates past VRAM and the device can write to it, so it looks like it works, but a paired A/B against pinned was identical (7.64/7.77 vs 7.63/7.77 tok/s) and `mem_info_gtt_used` stayed flat at 0.43 GiB while `vram_used` filled to 15.76 — consistent with the managed-memory note in `docs/ROCM.md` (`:56` on current main; the spike cited `:148` before that file was rewritten), so on discrete AMD the pinned-host path is what works and managed memory is not a shortcut to a UVA tier. `Backend::AllocManaged`/`FreeManaged` default to `nullptr` meaning "this backend has no managed allocator", so a caller falls back rather than assuming. **What the spike is NOT, and its code must not be lifted:** no `WeightOffloader`, no canonical-name targeting, no `cpu_offload_gb`, no `supports_weight_offload`; selection is raw byte size against a counter. Caveats stated rather than implied: one board, one ROCm version, one model family; the measurements sit on a base now **259 commits stale**, and both [#1402](https://github.com/mudler/vllm.cpp/pull/1402) and [#523](https://github.com/mudler/vllm.cpp/pull/523) landed afterwards and change decode cost, so the RATIOS are the durable part and the absolute tok/s figures are not; host contention was not controlled to benchmark standard; finding 3's mechanism is unexplained and findings 2 and 4 are single-board observations. Adjacent: llama.cpp's Vulkan backend loads the same 19.45 GiB file by spilling into the 31.35 GiB GTT the amdgpu driver exposes while its HIP backend refuses as we do ([#1400](https://github.com/mudler/vllm.cpp/issues/1400)), and [#1870](https://github.com/mudler/vllm.cpp/issues/1870) makes keep-quant residency load-bearing on a 16 GiB card, which changes what an offload budget competes for | record | +| [#1970](https://github.com/mudler/vllm.cpp/issues/1970) | `MODEL-DSV4-EXL3` | **The EXL3 loader asks for the DSA family at the COLLAPSED synthetic geometry, so the real DeepSeek-V4-Flash artifact shape-refuses on 41 of its 43 layers and every non-DSA capability behind it is unreachable.** Option C of the three [#1961](https://github.com/mudler/vllm.cpp/issues/1961) returned as `NEEDS_DECISION`, and a strict prefix of the full DSA port: the loader derives the DSA widths the way upstream derives them — STRICTLY, one width per layer, refusing anything else by name — and the forward REFUSES BY NAME instead of indexing a tensor at a width it does not have. The first cut accepted TWO widths (upstream's and a collapsed one) to keep a synthetic fixture loading; the fresh review showed that premise was not reproducible, because the four synthetic DSA suites contain zero references to `LoadDeepseekV4*` or `dsv4_exl3_fixture` and cannot break, and `coff` is a pure function of `compress_ratio` sizing `ape` (`:272`), both halves of `fused_wkv_wgate` (`:281`) and `state_cache.state_dim` (`:291`), so a `cr == 4` UNDOUBLED checkpoint is one upstream cannot load at all. Accepting it was a divergence from the mirror and is gone; the fixture moved to `cr == 128`, where `coff` is 1 and the collapsed width IS the derived one. NOT a new safety regression either way, and stated rather than implied: pre-PR (`git show c00625141:...deepseek_v4_weights.cpp`) the loader required exactly `{hd, H}`, so that malformed checkpoint was ALREADY accepted and ALREADY ran the collapsed `win = 2` maths — the derived form is the first version that refuses it. Four tensors refuse today, all on the 21 `compress_ratio == 4` layers — `attn.compressor.ape` `[4,1024]` against `[4,512]`, `attn.compressor.wgate.weight` `[1024,4096]` against `[512,4096]`, `attn.indexer.compressor.wkv.weight` `[256,4096]` against `[128,4096]`, and `attn.indexer.wq_b.weight` `[8192,1024]` against `[8192,4096]`. The doubled dimension is upstream's `coff = 1 + (compress_ratio == 4)` (`vllm/models/deepseek_v4/compressor.py:247-248` at the parity pin `5559679229bc961848b121ccdeaa8fa5d79bec98`), spent on the APE table (`:270-277`) and the fused projection (`:279-287`) and NOT on the norm (`:288` is `RMSNorm(self.head_dim, self.rms_norm_eps)`; `:293`, cited in this row before the fresh review, is `compress_ratio=compress_ratio` inside the `CompressorStateCache` call); the two halves are the two overlapping compression windows a token belongs to, selected at gather time by window position (`common/ops/fused_compress_quant_cache.py:164-183`) and not recoverable from the tensor alone. `indexer.wq_b` is not a width problem at all — upstream builds it as `ReplicatedLinear(q_lora_rank, head_dim * n_head)` (`attention.py:721-726`) and calls it on `qr` in `DeepseekV4Indexer.forward` (`:835`) while our forward feeds it the hidden state. The loader could NOT simply widen without the forward moving with it, and BOTH HALVES LAND TOGETHER — but the reason is DIAGNOSTIC and this row said otherwise before its fresh review. It claimed `Gemm`'s host arm is a `MatVec` with no length check, so that a `[1024,4096]` `comp_wgate` in a slot indexed as `[512,4096]` would be a silently wrong number. **That is false.** `deepseek_v4.cpp:413` is `VT_CHECK(w.size() == out * in, ...)`, unconditional, and `VT_CHECK` (`include/vt/dtype.h:11`) is a plain throw rather than an `assert`, so `NDEBUG` does not remove it; `Gemm` (`:428`) takes its keep-quant branch only when `be.gguf != nullptr` and an EXL3 load has `gguf == nullptr`, so the EXL3 DSA tensors take the checked unquantized arm, and the keep-quant arm checks too. NEITHER arm is unchecked. What the widened load without the refusal actually produces is an ANONYMOUS `vt: MatVec weight size mismatch at deepseek_v4.cpp:413` from the middle of a forward, naming no tensor, no layer, no geometry and nothing missing — verified by the fresh reviewer, who deleted the production call site while keeping the helper referenced so it compiled under `-Werror` and got that throw rather than logits. The refusal replaces an anonymous crash with a precise named refusal. It is a DIAGNOSTICS improvement, not the difference between wrong tokens and a refusal, and overstating it is the same class of false justification #1964 was filed for. `compress_ratios` was ALREADY read per layer and needed no change; only the widths derived from it were wrong. Explicitly NOT fixed here and owed on: the DSA maths itself (option A, no owning row), dense MLA as a fallback for `cr != 0` layers (that IS the [#1964](https://github.com/mudler/vllm.cpp/issues/1964) defect), the GGUF arm's `dsa_dense` behaviour (#1964, unchanged by the dispatch's own exclusion), the `cr == 128` EXL3 layers whose widths match while their `win = 2` pooling (`deepseek_v4.cpp:833`) is still not upstream's 128-wide boundary-emitted compressor over its own `compressor.wkv` projection — which is [#1976](https://github.com/mudler/vllm.cpp/issues/1976), filed by the fresh-review repair, and NOT #1964 as this row's spec first said, because #1964 is the GGUF arm's `dsa_dense` and closing it would not have closed this, and the `indexer.wq_b` input-space defect. Spec [`specs/dsv4-dsa-loader-accept-forward-refuse.md`](specs/dsv4-dsa-loader-accept-forward-refuse.md) | bug | | [#1960](https://github.com/mudler/vllm.cpp/issues/1960) | `KV-DSV4-MULTICACHE` | **`SlidingWindowMLASpec` is a declared enumerator with no struct behind it, and `MLAAttentionSpec` carries none of the four DeepSeek-V4 fields, so 105 of V4's 167 cache entries cannot be sized at all.** W1 of [#1925](https://github.com/mudler/vllm.cpp/issues/1925). `KVCacheSpecKind::kSlidingWindowMla` is declared at `include/vllm/v1/kv_cache_interface.h:89` and the port's deferral list names the class as omitted (`:46-52`); it is the spec class of the SWA cache (43 entries, `vllm/v1/attention/backends/mla/sparse_swa.py:86-101`) and of both compressor-state populations (41 + 21, `vllm/models/deepseek_v4/compressor.py:188-200`). `MLAAttentionSpec` (`kv_cache_interface.h:242-261`) adds no fields over `FullAttentionSpec` where upstream carries `cache_dtype_str`, `alignment`, `compress_ratio` and `model_version` (`vllm/v1/kv_cache_interface.py:381-388`), so the compressed latent is sized `block_size` rows per page where upstream stores `block_size // compress_ratio`, and the 584-byte `fp8_ds_mla` token (`:396-405`) throws by name instead (`src/vllm/v1/kv_cache_interface.cpp:64-71`). `_apply_alignment_padding` (`:345-351`) has no twin, so no V4 page reaches its 576B/512B alignment. Pure allocation metadata: nothing constructs either spec outside tests, because publishing before W3 would allocate a silent subset (`src/vllm/v1/worker/gpu/runner.cpp:577-597` drops an unmatched group kind with no diagnostic). | bug | +| [#1976](https://github.com/mudler/vllm.cpp/issues/1976) | `MODEL-DSV4-EXL3` | **The EXL3 arm's `cr == 128` DeepSeek-V4 layers run a 2-wide pool over the MLA's own latent where upstream runs a 128-wide boundary-emitted compressor over its own projection.** Split out of [#1970](https://github.com/mudler/vllm.cpp/issues/1970) during its fresh review, which found that [`specs/dsv4-dsa-loader-accept-forward-refuse.md`](specs/dsv4-dsa-loader-accept-forward-refuse.md) `## Owed` attributed this to [#1964](https://github.com/mudler/vllm.cpp/issues/1964). **That attribution is wrong and nothing else tracked it**, so closing #1964 would have closed a defect that is still live. The two are on DIFFERENT ARMS: #1964 is `dsa_dense = (be.gguf != nullptr)` (`src/vllm/model_executor/models/deepseek_v4.cpp:776`) making `is_comp` and `is_indexer` false on every layer, so a GGUF DeepSeek-V4 runs dense MLA where upstream runs the compressor and the "EXACT, not an approximation" justification beside it (`:758-770`) is false. This is the EXL3 arm, where `be.gguf` is null, `dsa_dense` is FALSE, and a `cr == 128` layer ENTERS the compressor: its widths already match, because `coff = 1 + (compress_ratio == 4)` is 1 at `cr == 128` (`vllm/models/deepseek_v4/compressor.py:247-248` at the parity pin `5559679229bc961848b121ccdeaa8fa5d79bec98`), so #1970's width refusal passes it through BY DESIGN and it then runs `const int64_t win = 2` (`deepseek_v4.cpp:833`) over the MLA's own `kraw` latent, emitted every token. Upstream instead pools a `coff * compress_ratio` = 128-wide window over a SEPARATE `compressor.wkv` projection and emits a row only at boundary tokens, `(position + 1) % compress_ratio == 0` (`compressor.py:171-173`), into a compressed KV cache distinct from the raw one. **Three things differ, not one**: window width (2 against 128), emission cadence (every token against every 128th) and source projection (the MLA latent against `compressor.wkv`, which #1970's loader accounts for and deliberately routes nowhere). This is why the real artifact's 20 `cr == 128` layers "already loaded before #1970" — they load, they run, and what they run is not upstream's compressor. Fix belongs to the DSA composition (option A of [`specs/dsv4-dsa-geometry.md`](specs/dsv4-dsa-geometry.md), [#1961](https://github.com/mudler/vllm.cpp/issues/1961)), which also needs the compressed-KV cache topology [#1960](https://github.com/mudler/vllm.cpp/issues/1960) and [#1925](https://github.com/mudler/vllm.cpp/issues/1925) are scoping. Filed separately rather than folded into #1961 because #1961 scopes the `coff == 2` overlapping-window pair while this layer class has `coff == 1` and passes every width check there is. NOT established and stated rather than implied: no token-level divergence has been MEASURED against the oracle for a `cr == 128` layer, because that needs the 99.5 GiB artifact and the box — the claim is a source-level one about window width, cadence and projection; no fix is proposed; and `CompressorSaveScoreApe` / `CompressorPoolNorm` are already generic over width and window, so the gap is the composition and the cache rather than the maths | bug | | [#1762](https://github.com/mudler/vllm.cpp/issues/1762) | `GEMMA4-FP8-WMMA-EXPERT-GEMM` | KEEP Gemma-4 FP8 T>1 expert path still dequantizes to BF16 and calls hipBLAS Tensile; a gated gfx1201 FP8 WMMA expert GEMM is the unblocked L2 lever (31.5% prefill). Spec-first, default-OFF, no GPU on this filing | perf | | [#526](https://github.com/mudler/vllm.cpp/issues/526) | `SERVE-TOOL-HISTORY-ARGS` | OpenAI multi-turn tool history reaches chat templates with string-valued arguments | bug | | [#1934](https://github.com/mudler/vllm.cpp/issues/1934) | `BACKEND-ROCM` | `RocmPlatform::needs_weight_staging()` is stale-false (a W0-era placeholder never revisited despite #523/#509/#506/ROCM_ATTN/hipGraph landing since), so `CheckDeviceWeightFit` — the #1123/#1870 load-time refusal, including the `policy_forces_full_expand` fix — never runs on ROCm: measured directly, `VT_DEVICE_WEIGHT_BUDGET_BYTES=1` produced no refusal on a real load. The actual device allocation the refusal guards is not gated on this flag, so #1870's crash stays reachable until this closes; owed, not fixed in flow, because flipping the flag also moves `DirectDeviceLoadEligible` and several GDN kernel-dispatch defaults that each need their own correctness check | bug | diff --git a/.agents/specs/dsv4-dsa-geometry.md b/.agents/specs/dsv4-dsa-geometry.md new file mode 100644 index 000000000..77813568c --- /dev/null +++ b/.agents/specs/dsv4-dsa-geometry.md @@ -0,0 +1,316 @@ +# DSV4-DSA-GEOMETRY — what the real DeepSeek-V4-Flash DSA tensors are, and why a reshape does not reach them + +Status: **SCOPING ONLY. No code lands from this document.** It answers the four +questions `MODEL-DSV4-EXL3` `## Owed` needed answered before the real 99.5 GiB +EXL3 artifact could load, and it reports that the answer is not the one that +entry assumed. The design choice that follows is a `NEEDS_DECISION` returned to +the operator, not a choice this document makes. + +Issue: [#1961](https://github.com/mudler/vllm.cpp/issues/1961) +Owning row: `MODEL-DSV4-EXL3` +Oracle: vLLM, primary, at the parity pin `5559679229bc961848b121ccdeaa8fa5d79bec98` +(`.agents/upstream-sync.md`), checked out at `/home/mudler/_git/vllm`. Every +`file:line` below is read at that pin. No secondary oracle is used or needed: +vLLM registers and implements this architecture in full. + +Local anchors: every `file:line` into THIS tree is read at the head of +`row/DSV4-DSA-GEOMETRY` as it lands, NOT at `c00625141` where the measurement +was taken. Seven of the eight distinct local citations below went stale INSIDE +this pull request, because #1970's implementation landed in +`deepseek_v4.cpp` and `deepseek_v4_weights.cpp` beside this document — the same +mechanism that put a wrong `attention.py` line into seven places on this row. +Where #1970 also changed the BEHAVIOUR a paragraph reports, the paragraph says +so instead of being re-pointed at a line that now reads the other way. + +## The measurement, taken first + +`compress_ratios` in the artifact's own `config.json` +(`/mnt/nas_share/rc/ckpt/dsv4-flash-0731-spark-exl3/config.json`) is a +**per-layer list**, not a scalar: `[0, 0, 4, 128, 4, 128, ..., 4, 0, 0, 0]` over +46 entries (43 layers + 3 MTP). Layers 0 and 1 are `0`; layers 2..42 alternate +`4`, `128`; the MTP tail is `0`. That is **21 layers at `cr == 4`**, **20 at +`cr == 128`**, **2 dense** — which is exactly the "41 of 43 carry a compressor, +21 carry an indexer" the row recorded, now with the reason attached. + +Read from the real shard headers (headers only, `carried-00{1..5}.safetensors`), +the DSA tensors split cleanly by that value: + +| tensor | `cr == 4` layers | `cr == 128` layers | our loader wants | +|---|---|---|---| +| `attn.compressor.ape` | `[4, 1024]` | `[128, 512]` | `{cr, 512}` | +| `attn.compressor.wgate.weight` | `[1024, 4096]` | `[512, 4096]` | `{512, 4096}` | +| `attn.compressor.wkv.weight` | `[1024, 4096]` | `[512, 4096]` | accounted only | +| `attn.indexer.compressor.ape` | `[4, 256]` | — | accounted only | +| `attn.indexer.compressor.wgate.weight` | `[256, 4096]` | — | accounted only | +| `attn.indexer.compressor.wkv.weight` | `[256, 4096]` | — | `{128, 4096}` | +| `attn.indexer.weights_proj.weight` | `[64, 4096]` | — | `{64, 4096}` ✓ | +| `attn.indexer.wq_b.weight` | `[8192, 1024]` | — | `{8192, 4096}` | + +So **four** tensors refuse, not the three `## Owed` names, and they refuse on the +21 `cr == 4` layers only. `compressor.ape` is the one the row had not counted: +`RequireShape` (`src/vllm/model_executor/models/deepseek_v4_weights.cpp:402-411`) +compared `{cr, hd}` = `{4, 512}` against the stored `[4, 1024]` — the loader as +it stood at `c00625141`, before #1970 made the expected width `{cr, coff * hd}`. +Every `cr == 128` +layer already satisfies our expectations byte for byte, and layers 0 and 1 carry +no compressor at all. + +## Q1 — what are the two halves of the doubled dimension? + +They are **the two overlapping compression windows a token belongs to**, not a +gate/value pair, not an interleave, and not a fusion. + +Upstream derives the width from one line: + +``` +vllm/models/deepseek_v4/compressor.py:247-248 + self.overlap = compress_ratio == 4 + self.coff = 1 + self.overlap +``` + +and spends it on the projection and the APE table: + +``` +vllm/models/deepseek_v4/compressor.py:279-287 + self.fused_wkv_wgate = MergedColumnParallelLinear( + self.hidden_size, + [self.coff * self.head_dim, self.coff * self.head_dim], # wkv | wgate + ... +vllm/models/deepseek_v4/compressor.py:270-277 + self.ape = nn.Parameter(torch.empty((compress_ratio, self.coff * self.head_dim), ...)) +``` + +`coff * head_dim` is `2 * 512 = 1024` for the main compressor at `cr == 4` and +`2 * 128 = 256` for the indexer's own compressor — every measured width above, +with no residue. At `cr == 128`, `overlap` is `False`, `coff` is `1`, and the +widths collapse to `head_dim`, which is why 20 of the 41 compressor layers load +today. + +`CompressorStateCache` states the same fact independently, and pins the domain: + +``` +vllm/models/deepseek_v4/compressor.py:171-173 + assert compress_ratio in [4, 128] + coff = 1 + (compress_ratio == 4) + self.sliding_window = coff * compress_ratio +``` + +The pooling window is `coff * compress_ratio` tokens wide while a compressed row +is emitted every `compress_ratio` tokens. At `cr == 4` that is an 8-token window +stepping 4, so **consecutive windows overlap by exactly half** and every token is +pooled twice — once as a member of the older half of a window, once as a member +of the newer half. The two halves of `wkv`/`wgate`/`ape` are the two projections +that serve those two roles. + +## Q2 — how does upstream index or split them? + +It never splits the weight. The split that `packed_modules_mapping` performs +(`vllm/models/deepseek_v4/nvidia/model.py:1157-1158`, and identically in +`amd/model.py:706-707`, `xpu/model.py:1166-1167`) is the *opposite* operation — +it **merges** the checkpoint's separate `compressor.wkv` and `compressor.wgate` +into one `MergedColumnParallelLinear` so the GEMM runs once. Our checkpoint +stores them unfused, which is the storage form upstream loads from. + +The `coff` halves are selected at **gather time, by window position**: + +``` +vllm/models/deepseek_v4/common/ops/fused_compress_quant_cache.py:164-183 + # in _fused_kv_compress_norm_rope_insert_sparse_attn (def at :114). Neither + # line below picks out its own line in the file: the indexer kernel repeats + # them at :712 and :730, the mxfp4-indexer kernel at :891 and :909. + if (position + 1) % COMPRESS_RATIO != 0: + return # boundary tokens only + start = position - (1 + OVERLAP) * COMPRESS_RATIO + 1 + tokens = tl.arange(0, (1 + OVERLAP) * COMPRESS_RATIO) # the coff*cr window + ... + head_offset = (tokens >= COMPRESS_RATIO).to(tl.int32) * HEAD_SIZE +``` + +`HEAD_SIZE` here is `head_dim` (512), not the stored width. A row in the older +half of the window (`tokens < cr`) reads its state at offset 0; a row in the +newer half reads at offset `head_dim`. The full `coff * head_dim` row is written +once per token by `save_partial_states` +(`vllm/models/deepseek_v4/common/ops/save_partial_states.py:86-101`, where +`HEAD_SIZE` is `kv.shape[-1]`, the *full* `coff * head_dim`, and the score half +gets `ape[position % cr]` added), and is read back twice, half at a time, by two +different windows. + +So the layout is a plain concatenation along the output dimension, addressed by a +role the token only acquires relative to the window doing the gathering. Nothing +about the halves is recoverable from the tensor alone. + +## Q3 — does our host forward want one half, both, or something else? + +**Neither.** Our host forward does not implement the composition these tensors +belong to, so there is no half to hand it. + +`AttentionBlock` (`src/vllm/model_executor/models/deepseek_v4.cpp:827-857`) +composes the compressor as: for **every** token `t`, softmax-pool a **fixed +`win = 2`** window of the **MLA's own `kraw` latent** and overwrite `latent[t]` +in place. The file says so itself at `:32-33`. Upstream, at the same point, +emits a row **only at `(position + 1) % cr == 0`**, pools **`coff * cr` = 8 or +128** rows, of a **separate `compressor.wkv` projection** that our loader +deliberately does not materialize, into a **separate compressed KV cache** that +sits beside the raw one. + +The differences are not parameters of one algorithm: + +| axis | upstream | ours | +|---|---|---| +| compressor KV source | `compressor.wkv`, its own projection | reuses MLA `kraw` | +| window width | `coff * cr` (8 or 128) | fixed 2 | +| emission | boundary tokens only, 1 row per `cr` | every token | +| destination | separate compressed cache | overwrites the dense latent | +| `coff` half selection | `head_offset` by window position | absent | +| indexer `wq_b` input | `qr`, the q-LoRA latent (`DeepseekV4Indexer.forward`, `attention.py:835`) | `x`, the hidden state (`deepseek_v4.cpp:915`) | +| indexer K | `indexer.compressor`, a pooled compressor | plain `Gemm(idx_wk, x)` (`:917`) | +| indexer selects among | compressed rows | raw rows | +| attention keys | SWA(128) raw ∪ selected compressed, one softmax | all raw rows, dense causal | + +The `wq_b` row is worth separating from the rest, because it is **not a width +problem at all** and it is a defect independent of this decision. Upstream builds +it as `ReplicatedLinear(self.q_lora_rank, self.head_dim * self.n_head)` +(`vllm/models/deepseek_v4/attention.py:721-726`) and calls it on `qr` in +`DeepseekV4Indexer.forward` (`:835`). The stored `[8192, 1024]` is therefore `[inh*ihd, q_lora_rank]` at its +natural size — `64*128` by `q_lora_rank == 1024` — and nothing about it is +doubled. At `c00625141` our loader asked for `[inh*ihd, H]` = `[8192, 4096]`; +#1970 has since moved that K to `q_lora_rank` +(`deepseek_v4_weights.cpp:995`), so the LOADER half is repaired. Our forward +still feeds it `x` (`deepseek_v4.cpp:915`), so we still project the indexer +query from the wrong space. + +No gate has ever seen that, and the reason is NOT that `H` and `q_lora_rank` +coincide at the synthetic geometry. They do not: `dsv4_exl3_fixture.h:141` +sets `kHidden` to 256 and `:149` sets `kQLora` to 128. The reason is that the +collapsed fixture WRITES `wq_b` at `K = H`, to match what our forward feeds it +— so the two agree by construction and the disagreement never appears. It +takes a fixture that writes the real geometry everywhere else and collapses +this one tensor to make it visible, which is what +`FixtureOptions::collapsed_indexer_wq_b` now does. + +## Q4 — is `dsa_dense` a mode upstream has? + +Upstream has a per-layer dense mode, and it is **config-driven, not +source-driven**: + +``` +vllm/models/deepseek_v4/attention.py:209 self.compress_ratio = max(1, config.compress_ratios[layer_id]) +vllm/models/deepseek_v4/attention.py:334 if self.compress_ratio > 1: # compressor exists +vllm/models/deepseek_v4/attention.py:274 if self.compress_ratio == 4: # indexer exists +vllm/models/deepseek_v4/nvidia/flashinfer_sparse.py:263 swa_only = self.compress_ratio <= 1 + # DeepseekV4FlashInferMLAAttention.forward_mqa; the same line is at :686 + # and :793 on DeepseekV4FlashInferSM120Attention. +``` + +`dsa_dense = (be.gguf != nullptr)` (`deepseek_v4.cpp:776`) keys off the +**weight source**. Upstream keys off `compress_ratios[layer_id]`. The two agree +on nothing: on this checkpoint the config-driven predicate is true for 2 layers +and false for 41, while ours is true for all 43 whenever the source is GGUF. So +`dsa_dense` is our workaround, and widening it to the EXL3 source is wrong for +the reason the task states and for a second one below. + +**And upstream's dense layers are not dense.** `swa_only` means *sliding-window +only*, at `config.sliding_window == 128` (`attention.py:204`). Even layers 0 and +1 attend a 128-token window, never the full prefix. Our forward has no sliding +window anywhere (`grep -n sliding_window src/.../deepseek_v4.cpp` returns only +prose). + +## The claim `dsa_dense` rests on is false, and that is the finding + +`deepseek_v4.cpp:763-775` justifies forcing the DSA path off with: dense MLA "is +EXACT, not an approximation, whenever `seq_len <= index_topk` (=512): the indexer +cannot select more tokens than exist, so top-k over ≤512 tokens IS the full +causal set". [#1925](https://github.com/mudler/vllm.cpp/issues/1925) repeats it. + +The premise is right about the indexer and wrong about the attention. On a +`cr > 1` layer, one kernel takes **one softmax over the union** of the raw +sliding window and the selected compressed rows: + +``` +vllm/models/deepseek_v4/nvidia/flashinfer_sparse.py:769-782 + # DeepseekV4FlashInferSM120Attention._forward_decode; the same call is at + # :486, :511 and :888. + flashinfer_trtllm_batch_decode_sparse_mla_dsv4( + query=q, + swa_kv_cache=swa_cache, sparse_indices=swa_indices, + compressed_kv_cache=extra_cache, extra_sparse_indices=extra_sparse_indices, + sinks=self.attn_sink, ...) +``` + +The compressed rows are **pooled aggregates of `coff * cr` raw rows**, not raw +rows, so no selection over them can reproduce attention over raw rows. Their +number does not depend on `index_topk` either: rows are emitted at every +`(position + 1) % cr == 0`, so a 10-token prefill at `cr == 4` already has two, +and the short-context branch (`attention.py:813-830`) explicitly still builds the +K cache and selects **all** candidates — its comment says "we still need to build +k cache". Upstream at `seq_len == 10, cr == 4` attends 10 raw keys **and** 2 +compressed keys. We attend 10. + +Upstream's attention here is hierarchical: recent tokens at full resolution, +older tokens pooled `cr:1`, jointly normalized. Dense causal attention is not +that at any sequence length. **A token gate below 512 cannot detect this, and +neither can one above it.** The consequence reaches past this row: the GGUF arm +runs `dsa_dense` on the real geometry today, so the shipping GGUF DeepSeek-V4 +path is already not upstream's attention on 41 of 43 layers. + +## Why this is `NEEDS_DECISION` and not a patch + +The loader-side half is small and unambiguous — derive the expected widths as +upstream does (`coff = 1 + (cr == 4)`) and take `wq_b`'s K from `q_lora_rank`. +But landing only that is worse than the refusal it removes, and #1970 corrected +what "worse" means here. This paragraph used to say `Gemm`'s host arm is a +`MatVec` with no length check, so the wrong stride would be read SILENTLY. It is +not: `deepseek_v4.cpp:413` is an unconditional `VT_CHECK` and the keep-quant arm +checks the shape too. The moment `comp_wgate` materializes at `[1024, 4096]` and +`AttentionBlock` calls `Gemm(..., T, hd, H)` with `hd == 512`, what happens is an +ANONYMOUS `vt: MatVec weight size mismatch` from the middle of a forward, on a +checkpoint that loaded successfully, naming no tensor, no layer and nothing +missing. That is a worse DIAGNOSTIC than the loader refusal it replaced, not a +worse numerical outcome, so the refusal must not move without the forward moving +with it. + +Three shapes the decision can take. This document recommends none. + +- **(A) Port upstream's DSA.** Separate `compressor.wkv`; the `coff`-overlapped + window with `head_offset` role selection; boundary emission; a compressed KV + cache beside a SWA(128) raw cache; the indexer on `qr` over compressed rows; + one joint softmax over the union. The only path to token parity with vLLM on + this checkpoint. It is a multi-wave model port and it needs the cache topology + [#1960](https://github.com/mudler/vllm.cpp/issues/1960) and + [#1925](https://github.com/mudler/vllm.cpp/issues/1925) are already scoping — + 105 of V4's 167 cache entries cannot be sized at all today. Our primitives + (`CompressorSaveScoreApe`, `CompressorPoolNorm`) are already generic over + width and window and would carry over unchanged; the gap is the composition, + not the math. +- **(B) A per-layer dense selector both arms read**, mirroring + `compress_ratios[layer_id]` rather than the weight source. Honest as a record + and it fixes the arm-divergence the row's equivalence gate would otherwise + suffer, but on 41 layers it is still not upstream's attention, so it cannot be + gated as parity and must not be described as exact. +- **(C) Loader accepts, forward refuses by name.** Materialize every DSA tensor + at its real width and move the refusal into `AttentionBlock`. The artifact + then loads and its non-DSA capabilities — the EXL3 tower at real scale, MoE, + MTP, W2 residency — become reachable, with the DSA layers refusing instead of + mis-indexing. Under `## Nothing lands dead` this is a staged slice and owes a + named owner for the wiring. + +(A) is what "mirror vLLM" means here. (C) is what unblocks the row this week. +They are not exclusive — (C) is a strict prefix of (A) — but which one is dispatched, +and what the row's equivalence gate compares against once the two arms stop +sharing an attention path, exceeds a helper's authority. + +## Owed + +- The decision above. +- `MODEL-DSV4-EXL3` `## Owed` names three refusing tensors; there are four. + `attn.compressor.ape` is missing from it, and `indexer.wq_b` is listed there + as a width problem when it is a wrong-input-space problem. +- The `indexer.wq_b` input-space defect (`x` where upstream uses `qr`) is real at + any geometry and is not fixed by any of (A)/(B)/(C) on its own. +- The GGUF arm's `dsa_dense` exactness claim (`deepseek_v4.cpp:763-775`) is + false and is quoted onward by #1925. Whoever takes the decision owns + correcting both prose sites. + +## Now + +`SCOPING`. No lifecycle state moved. No product code changed. diff --git a/.agents/specs/dsv4-dsa-loader-accept-forward-refuse.md b/.agents/specs/dsv4-dsa-loader-accept-forward-refuse.md new file mode 100644 index 000000000..2a032b8ad --- /dev/null +++ b/.agents/specs/dsv4-dsa-loader-accept-forward-refuse.md @@ -0,0 +1,672 @@ +# DSV4-DSA-ACCEPT-REFUSE — the loader takes the real DSA geometry, the forward refuses by name + +Issue: [#1970](https://github.com/mudler/vllm.cpp/issues/1970) +Owning row: `MODEL-DSV4-EXL3` +Scoped by: [`.agents/specs/dsv4-dsa-geometry.md`](dsv4-dsa-geometry.md) ([#1961](https://github.com/mudler/vllm.cpp/issues/1961)) +Oracle: vLLM, primary, at the parity pin `5559679229bc961848b121ccdeaa8fa5d79bec98` +(`.agents/upstream-sync.md`), read at `/home/mudler/_git/vllm`. Every `file:line` +below is read at that pin. No secondary oracle is used or needed: vLLM registers +and implements this architecture in full. + +This is **option C** of the three the geometry spec returned as `NEEDS_DECISION`. +It is a **strict prefix of option A** (the full DSA port) and forecloses nothing. + +## Scope + +1. The EXL3 loader materializes every DSA tensor at the width the REAL + DeepSeek-V4-Flash artifact stores, derived the way upstream derives it, so the + real checkpoint loads instead of shape-refusing on 41 of 43 layers. +2. `AttentionBlock` verifies, before it indexes any DSA tensor, that the + materialized width is the one its own arithmetic assumes, and REFUSES BY NAME + when it is not. + +Nothing else moves. In particular the DSA maths is not ported, dense MLA does not +become a fallback, and the GGUF arm's behaviour is unchanged. + +## Upstream anchors + +The whole geometry follows from one line: + +``` +vllm/models/deepseek_v4/compressor.py:247-248 + self.overlap = compress_ratio == 4 + self.coff = 1 + self.overlap +``` + +and is spent on exactly three parameters, plus a norm that is NOT widened: + +| upstream | anchor | width | +|---|---|---| +| `self.ape` | `compressor.py:270-277` | `[compress_ratio, coff * head_dim]` | +| `fused_wkv_wgate` | `compressor.py:279-287` | two outputs of `coff * head_dim` | +| `self.norm` | `compressor.py:288` | `RMSNorm(self.head_dim, self.rms_norm_eps)` — **not** `coff * head_dim` | +| `indexer.wq_b` | `attention.py:721-726` | `ReplicatedLinear(q_lora_rank, head_dim * n_head)` | + +The indexer carries its OWN `DeepseekCompressor` at `head_dim = index_head_dim` +and the SAME `compress_ratio` (`attention.py:768-776`), so its `ape`/`wgate`/`wkv` +are `coff * index_head_dim` wide and its `norm` is `index_head_dim` wide. The +indexer exists only at `compress_ratio == 4` (`attention.py:274`, the `if` +itself; `:276` is a comment inside it), where `coff` is 2, which is why only +`cr == 4` layers refuse today. + +Every anchor above was re-verified against the checkout at the pin during the +fresh-review repair, asserting UNIQUENESS and not mere existence. `:288` is the +only `RMSNorm(` in `compressor.py`; `:293`, cited here before the repair, is +`compress_ratio=compress_ratio` inside the `CompressorStateCache` call. + +`compress_ratio` is per layer upstream — `max(1, config.compress_ratios[layer_id])` +(`attention.py:209`) — and our `DeepseekV4Params::compress_ratio(layer)` +(`include/vllm/model_executor/models/deepseek_v4.h:122-128`) already mirrors that, +including `has_indexer == (cr == 4)`. **No config-parsing change is owed**; the +per-layer list was already read correctly. The defect was only in the widths the +loader derived from it. +## Design + +### D1. The loader DERIVES the width the way upstream derives it + +`deepseek_v4_weights.cpp` gains upstream's own expression, + +``` +const int64_t coff = (cr == 4) ? 2 : 1; // compressor.py:247-248 +``` + +and a `RequireDsaDim` helper that refuses any width but the derived one BY NAME. +Three DIFFERENT rules produce the three widths, so the helper takes a `why` string +and each call site names its own derivation: + +| tensor | required | derivation | +|---|---|---| +| `compressor.ape`, `compressor.wgate.weight` | `coff * head_dim` | `compressor.py:247-248`, `:270-277`, `:279-287` | +| `indexer.compressor.wkv.weight` | `2 * index_head_dim` | the indexer's own compressor at the same ratio (`attention.py:768-776`), which exists only at `cr == 4` (`:274`) | +| `indexer.wq_b` dim 1 | `q_lora_rank` | `ReplicatedLinear(q_lora_rank, head_dim * n_head)` (`attention.py:721-726`) called on `qr` in `DeepseekV4Indexer.forward` (`:835`) — **not** a `coff` width, and the refusal must not cite one | + +`compressor.norm.weight` stays `{hd}` and `indexer.weights_proj` stays `{inh, H}`, +because upstream widens neither. + +**Why it derives ONE width and does not accept two.** `coff` is a pure function of +`compress_ratio`. It sizes `ape` (`:272`), both halves of `fused_wkv_wgate` +(`:281`) and `state_cache.state_dim` (`:291`); the indexer is gated on +`compress_ratio == 4` (`attention.py:274`) and builds its own compressor at the +same ratio (`:768-776`). So for a given ratio upstream emits exactly ONE width, +and a `cr == 4` checkpoint carrying an UNDOUBLED family is one upstream cannot +load at all. `AGENTS.md` requires this loader to mirror every mode, default, error +and edge case upstream defines, so accepting a width upstream can never emit is a +divergence from the mirror — and production acceptance must not be widened to +suit a fixture. + +The first cut of this row DID accept two widths, on the premise that refusing the +collapsed one would delete the synthetic gates. **That premise was not +reproducible, and the fresh review is what caught it.** +`test_deepseek_v4_compressor.cpp`, `test_deepseek_v4_dsa.cpp`, +`test_deepseek_v4_forward.cpp` and `test_deepseek_v4_mtp.cpp` contain ZERO +references to `LoadDeepseekV4*` or `dsv4_exl3_fixture`, so they never load through +this arm and cannot break. Only the two EXL3 suites did, and this change rewrites +their fixture anyway: `ForwardFixtureOptions()` moves from `compress_ratios = +{0, 4}` to `{0, 128}`, where `coff` is 1 and the derived width IS the collapsed +one, so the W2d, MoE, #1923 and residency cases keep driving an EXL3-loaded +compressor layer end to end through the production forward. + +**What that costs, stated rather than implied.** No case runs an EXL3-loaded +INDEXER forward any more, because the indexer exists only at `cr == 4`, where the +real geometry is the one the forward refuses. That coverage is redundant with +`test_deepseek_v4_forward.cpp` and `test_deepseek_v4_dsa.cpp`, which exercise the +indexer maths at the synthetic geometry without loading through this arm, and the +real artifact refuses it in any case. The indexer's LOAD is still gated at +upstream's widths by the loader suite's tower case, which moves to +`real_dsa_geometry = true`. + +**The derived form is also strictly safer than what stood before this row.** +Pre-PR (`git show c00625141:src/vllm/model_executor/models/deepseek_v4_weights.cpp`) +the loader required exactly `{hd, H}`, so a malformed `cr == 4` UNDOUBLED +checkpoint was ALREADY accepted and ALREADY ran the collapsed `win = 2` maths. +The two-width form was therefore not a new safety regression, and the derived form +improves on both — it is the first version that refuses that checkpoint. + +**Half-widening needs no separate rule, and this is a structural claim rather +than a gated one — stated that way because the version before the fresh-review +repair asserted it with no gate at all.** Every member is checked independently +against the same derived value: `ape` through `Float`'s `RequireShape` at +`{cr, cw}`, `wgate` and the indexer's `wkv` through `RequireDsaDim`, `wq_b` +through `RequireDsaDim` on dim 1. There is no cross-member "family agrees with +itself" logic left that could be wrong, so a checkpoint that widens `wgate` but +not `ape` reds on `ape`. **What is gated is the fully collapsed family (M8) and +`indexer.wq_b` alone (M11); a MIXED half-widened fixture is not gated, and no +knob writes one.** M11 exists because it must: the compressor refuses before the +loader reaches `wq_b`, so without `collapsed_indexer_wq_b` that derivation would +be unfalsifiable no matter how many collapsed cases were added. + +At `cr == 128`, `coff` is 1, the derived width is the collapsed one, and the 20 +`cr == 128` layers that load today keep loading unchanged. + +### D2. The forward checks its own preconditions and refuses + +**What this is worth, stated exactly, because the first cut of this row +overstated it and the fresh review rejected the claim.** `Gemm`'s host arm is a +`MatVec` whose size assertion is **unconditional**: `deepseek_v4.cpp:413` is +`VT_CHECK(w.size() == out * in, "MatVec weight size mismatch")`, and `VT_CHECK` +(`include/vt/dtype.h:11`) is a plain `throw`, not an `assert`, so `NDEBUG` does +not remove it. That is also the arm the EXL3 DSA tensors take, because `Gemm` +(`:428`) enters its keep-quant branch only when `be.gguf != nullptr` and an EXL3 +load has `be.gguf == nullptr` — and the keep-quant branch checks the shape too. +NEITHER arm is unchecked. + +So a wide `comp_wgate` read at a `[hd, H]` stride does **not** produce a plausible +wrong number. It produces + +``` +vt: MatVec weight size mismatch at deepseek_v4.cpp:413 +``` + +from the middle of a forward, on a checkpoint that loaded successfully, naming no +tensor, no layer, no geometry and nothing missing. **This is therefore a +DIAGNOSTICS improvement and that is the whole of it** — an anonymous crash +replaced by a precise named refusal. It is not the difference between wrong tokens +and a refusal. Three review rounds were needed to make every reachable copy say +so, and each round reported the sweep complete before the next one found more. +Round 3 found a TWELFTH, in `deepseek_v4_weights.cpp`'s carried-half block: it +read "names the tensor rather than producing a wrong number" in the same file +whose reader-shape paragraph already says the opposite. It is repaired. One copy +is beyond repair and is named here rather than reported as swept: the +[#1923](https://github.com/mudler/vllm.cpp/issues/1923) row in +`.agents/issue-index.md`, which is on `origin/main` and append-only. +Overstating it is exactly the class of false justification +[#1964](https://github.com/mudler/vllm.cpp/issues/1964) was filed for, and this +row must not repeat it one directory over. + +`AttentionBlock` therefore checks, for every DSA tensor it is about to index, +that the materialized element count equals the count its indexing assumes: + +| slot | indexed as | anchor | +|---|---|---| +| `comp_ape` | `[cr, hd]` | `DispSaveScoreApe(..., T, hd, cr)` | +| `comp_wgate` | `[hd, H]` | `Gemm(..., T, hd, H)` | +| `comp_norm_weight` | `[hd]` | `DispPoolNorm(..., hd)` | +| `idx_wq` | `[inh * ihd, H]` | `Gemm(..., T, inh * ihd, H)` | +| `idx_wk` | `[ihd, H]` | `Gemm(..., T, ihd, H)` | +| `idx_wproj` | `[inh, H]` | `Gemm(..., T, inh, H)` | + +All six are checked, not only the four the real artifact stores differently. + +**Two of the six are not reachable from any checkpoint, and this says so rather +than leaving "all six are load-bearing" as an impression.** `comp_norm_weight` and +`idx_wproj` are the two slots upstream does not widen, so D1 requires them at +exactly the width this forward indexes and no artifact can make them disagree. +They are gated for FALSIFIABILITY only, by mutating a tower the production loader +produced (M9 and M10): that proves the checks fire and are named in the message, +and it proves nothing about reachability, because the state it constructs is one +the loader cannot emit. They are defensive checks against a future loader/forward +disagreement. The other four are what fires on the real checkpoint. + +The refusal names the layer, the tensor, both element counts, the width the +forward composes with, the width the checkpoint carries, the `coff` that explains +it, the missing capability, the upstream anchors and the issue. It is a +`VT_CHECK`, so it surfaces as the same refusal every other unrepresentable input +on this path surfaces as. + +**The check is gated on `is_comp || is_indexer`.** A layer whose DSA path the +forward does not enter reads none of these tensors, so checking it would refuse a +load that harms nothing. This is also what keeps the GGUF arm — where +`dsa_dense` makes both predicates false — byte-for-byte unchanged. + +### D3. What this deliberately does NOT decide + +At `cr == 128` the widths already match, so a `cr == 128` layer passes D2 and runs +the existing `win = 2` pooling of the MLA's own `kraw`. That is **not** upstream's +128-wide boundary-emitted compressor over a separate `compressor.wkv` projection, +and this change does not make it one. **It is +[#1976](https://github.com/mudler/vllm.cpp/issues/1976), filed by this repair, and +it is NOT #1964.** #1964 is the GGUF arm's `dsa_dense` routing every layer to +dense MLA; this is the EXL3 arm, where `dsa_dense` is false, the layer ENTERS the +compressor, and what it runs is a 2-wide pool over the MLA's own latent +(`deepseek_v4.cpp:833`) against upstream's 128-wide boundary-emitted compression +over a separate `compressor.wkv` projection (`compressor.py:171-173`). Window +width, emission cadence and source projection all differ. Closing #1964 would not +close this, and before #1976 nothing else tracked it. Recorded under `## Owed` rather than silently widened, because +widening D2 into "refuse every compressor layer" would refuse the gated synthetic +suites too, which is a scope change and not this unit of work. +## Risks + +- **The refusal could be unreachable.** The failure `.agents/reachability.md` + documents, and the exact failure #1923 already cost this row once: W2's + reachability claim was gated on a `DeepseekV4Weights` the suite built BY HAND, + so a load could not produce it. Mitigated by driving the RED test through + `vllm::LoadDeepseekV4ForCausalLMWeights` and `vllm::DeepseekV4Model::Forward`, + and by the reachability mutation in `## Gates` below. +- **The loader could widen without the forward moving.** Mitigated by landing + both halves in one change and by the reachability mutation `M6`, which deletes + the production call site. What that mutation observes is + `vt: MatVec weight size mismatch at deepseek_v4.cpp:413` — a THROW, and not a + wrong number. §D2 forty lines above says why: `MatVec`'s size assertion is + unconditional and `Gemm`'s keep-quant arm checks the shape too, so neither arm + is unchecked. The consequence of this risk is therefore an ANONYMOUS crash and + not wrong tokens. An earlier draft of this bullet said "a wrong number rather + than a throw"; that was the #1964 overclaim surviving in a seventh place, and + it contradicted this document's own §D2. +- **The fixture could describe a geometry the artifact does not have.** It + already did: `real_compressor_width` doubled the width unconditionally, so the + one existing case that used it wrote a doubled compressor on a `cr == 128` + layer, where upstream's `coff` is 1. Mitigated by replacing the flag with + upstream's own per-layer rule, so the fixture cannot describe a width upstream + would not produce. +## Tests + +`tests/vllm/models/dsv4_exl3_fixture.h` — `real_compressor_width` becomes +`real_dsa_geometry`, applying `coff = 1 + (cr == 4)` per layer to the compressor +and indexer-compressor families, leaving both norms at `head_dim`, and writing +`indexer.wq_b` at `[inh * ihd, q_lora_rank]`. `collapsed_indexer_wq_b` collapses +that ONE tensor while the rest of the family stays real, `collapsed_indexer_wkv` +does the same for `indexer.compressor.wkv.weight`, and `bogus_dsa_width` writes a +third width no oracle emits. + +`tests/vllm/models/test_deepseek_v4_exl3_forward.cpp` — `ForwardFixtureOptions()` +moves to `compress_ratios = {0, 128}`, where `coff` is 1 and the derived width is +the collapsed one, so every pre-existing case keeps running an EXL3-loaded +compressor layer end to end. `RealDsaFixtureOptions()` is the same model at +`{0, 4}` with the real geometry. One new case, entering through the production +loader and the production forward: + +1. The real geometry **LOADS**. Red before D1 (`RequireShape` throws). +2. The forward **REFUSES BY NAME**, and the message carries the layer, every + mismatched tensor, both counts and the missing capability. Red before D2 — and + red for the DIAGNOSTIC, not for the throw: without D2 the forward throws the + anonymous `MatVec weight size mismatch` from `deepseek_v4.cpp:413` instead, so + `!msg.empty()` alone does not separate the two and every `Mentions` assertion + after it is part of the gate. +3. A compressor layer the forward CAN index still runs. +4. `comp_norm_weight` and `idx_wproj` are gated by mutating the loaded tower, + which is a falsifiability gate and not a reachability one (D2). They are + mutated ONE PER FORWARD: at this fixture `head_dim - 1` and + `index_n_heads * hidden_size - 1` are the same number (511), so a single + message cannot say which tensor reported which count. + +Every count assertion in that case is bound to the mismatch LINE that names its +tensor, not asserted as a bare number. Several counts coincide here — +`compress_ratio * head_dim`, `index_n_heads * index_head_dim * hidden_size` and +`2 * index_head_dim * hidden_size` are all 2048 — so a bare `Mentions(msg, +"2048")` is satisfied by a tensor other than the one it was written for, which +the round-2 review measured as two more surviving mutants. `indexer.wq_b` gains +the counts it never had. + +`tests/vllm/models/test_deepseek_v4_exl3_loader.cpp` — the tower case moves to +`real_dsa_geometry = true` and asserts the doubled and the undoubled members side +by side, because an assertion that checked only the four doubled ones would pass a +loader that widened the whole family. Two new refusal subcases: a COLLAPSED +`cr == 4` family, which is the shape a two-width loader accepted and upstream +cannot emit; `indexer.wq_b` at `K = hidden_size` with the rest of the family +real, which is the only way to reach that check because the compressor refuses +first; and `indexer.compressor.wkv.weight` at the COLLAPSED `index_head_dim` +with the rest of the family real, which is the only way to reach the THIRD +derivation for the same reason. The round-2 fresh review found that third hole +by mutation: deleting the loader's `coff * index_head_dim` check left both suites +green, because no fixture wrote a real-geometry compressor beside a collapsed +indexer and the derivation's message was never read. The old compressor subcase stays retired: it asserted the behaviour this +change removes, and its `cr == 128` fixture doubled a layer where `coff` is 1. + +## Gates + +```sh +cmake -S . -B build -G Ninja -DVLLM_CPP_SERVER=OFF +cmake --build build --target test_deepseek_v4_exl3_forward test_deepseek_v4_exl3_loader -j 4 +ctest --test-dir build -R 'deepseek_v4' --output-on-failure +scripts/agent-preflight.sh +``` + +Build ONLY those targets. A bare `ninja -C build` links every test binary in the +tree, which took `build/tests` to 9.4 GiB during the fresh-review repair, took the +host to 100% full, and failed the link with `No space left on device`. Per +`.agents/environment.md` an ENOSPC here makes checkers emit FALSE policy refusals +rather than clean failures, so it is worse than a slow build. + +Every build records ninja's exit code AND its step count: a failed build silently +re-runs a stale binary and reads as a pass. + +Mutations, each verified to have LANDED before the build, then applied to a +scratch copy and restored byte-for-byte under SHA-256: + +| # | mutation | must | +|---|---|---| +| M1 | delete the `comp_wgate` width check | RED | +| M2 | delete the `comp_ape` width check | RED | +| M3 | delete the `idx_wq` width check | RED | +| M4 | delete the `idx_wk` width check | RED | +| M5 | loader reverts to the collapsed CONSTANT width (pre-#1970) | RED | +| M6 | **reachability**: delete the production call site in `AttentionBlock` | RED | +| M7 | `RequireDsaDim` accepts ANY width | RED | +| M8 | loader ALSO accepts the collapsed width (the two-width form D1 removes) | RED | +| M9 | delete the `comp_norm_weight` check | RED | +| M10 | delete the `idx_wproj` check | RED | +| M11 | `indexer.wq_b`'s K also accepts `hidden_size` | RED | +| M12 | `compressor.ape`'s EXPECTED count is wrong on the `cr == 4` layer | RED | +| M13 | `compressor.norm.weight`'s EXPECTED count is wrong on that layer | RED | +| M14 | delete the loader's `coff * index_head_dim` derivation for `indexer.compressor.wkv.weight` | RED | + +M8 through M11 were added by the round-1 repair, and three of them are that +review's findings in executable form. M9 and M10 were GREEN against the first cut, +which is what showed two of the six forward checks were unfalsifiable. M11 was +GREEN even after D1 became strict, because the compressor refuses before the +loader reaches `wq_b` — which is why `collapsed_indexer_wq_b` exists. + +M12 through M14 were added by the ROUND-2 repair and all three were GREEN against +`0acf0147f`. M14 is the loader's third derivation, unfalsifiable for exactly the +reason M11 was, one tensor over; `collapsed_indexer_wkv` is its +`collapsed_indexer_wq_b`. M12 and M13 are the count collisions: they leave the +tensor NAMED and make only its number wrong, which a bare `Mentions` of that +number cannot see because another tensor's line still carries it. + +**Three derivations, and now three of them falsifiable.** `RequireDsaDim` is +called from three sites deriving three DIFFERENT rules — `coff * head_dim`, +`coff * index_head_dim`, and `q_lora_rank`, which is not a `coff` width at all — +and until M14 the middle one had no case that read its message. M7 mutates +`RequireDsaDim` itself and so cannot tell the three apart. + +## Evidence + +Measured on `row/DSV4-DSA-GEOMETRY`, host build +`cmake -S . -B build -G Ninja -DVLLM_CPP_SERVER=OFF`, GCC with `-Wall -Wextra +-Werror`. These are the FRESH-REVIEW REPAIR's numbers, re-derived on the repaired +tree. The first cut's table is superseded and not carried forward, because D1 +changed shape and two of its mutations scored against a rule that no longer +exists. + +**RED first, D1.** The strict rule's own gate — the loader must refuse a COLLAPSED +`cr == 4` family — written before the loader changed, against the two-width form. +`ninja rc=0, 2/2 steps`, so the red is the test and not a stale binary. The +`:661` inside the transcript below is the line of the tree it was CAPTURED on +and is NOT a live anchor: at this branch's head that line is a comment. It is +labelled here because round 3 found it labelled only in the pull-request body: + +``` +tests/vllm/models/test_deepseek_v4_exl3_loader.cpp:661: ERROR: CHECK( Mentions(msg, "coff") ) is NOT correct! + values: CHECK( false ) + logged: msg := + +[doctest] test cases: 1 | 0 passed | 1 failed | 10 skipped +[doctest] assertions: 15 | 8 passed | 7 failed | +``` + +`msg` is EMPTY. The two-width loader accepted the collapsed `cr == 4` family and +threw nothing at all, which is the divergence from upstream this repair removes. + +**RED first, D2**, from the first cut and still valid — the new forward case +failing because the loader refused the real width before the forward could see it: + +``` +TEST CASE: dsv4 exl3 #1970: the REAL DSA geometry LOADS and the FORWARD refuses by name +ERROR: test case THREW exception: vt: deepseek-v4 exl3 loader: + layers.1.attn.compressor.ape must be [4,512], got [4,1024]. +[doctest] test cases: 1 | 0 passed | 1 failed | 4 skipped +[doctest] assertions: 0 | 0 passed | 0 failed | +``` + +That build was `ninja rc=0`, 500/500 steps. + +**A previously recorded D2 red was WRONG and is withdrawn.** This section used to +read "no throw: the forward returns logits computed off a mis-indexed +`comp_wgate`". It does not and cannot: `MatVec`'s size assertion is unconditional +(D2), so without the refusal the forward throws +`vt: MatVec weight size mismatch at deepseek_v4.cpp:413`. The fresh reviewer +demonstrated exactly that by deleting the production call site while keeping the +helper referenced so it compiled under `-Werror`, and got the throw rather than +logits. What D2 buys is the DIAGNOSTIC, and no record here may claim more. + +**RED first, round 2.** The three checks a mutation walked through at +`0acf0147f`, each measured GREEN there and RED after the repair. All six builds +were `ninja rc=0`. Verbatim, from the repaired tree: + +``` +test_deepseek_v4_exl3_loader.cpp:746: ERROR: CHECK( Mentions(msg, "coff * index_head_dim") ) is NOT correct! + values: CHECK( false ) +test_deepseek_v4_exl3_loader.cpp:754: ERROR: CHECK( Mentions(msg, "dimension 0 must be 8") ) is NOT correct! + values: CHECK( false ) +``` + +``` +test_deepseek_v4_exl3_forward.cpp:478: ERROR: CHECK( dsv4_exl3_fixture::Mentions( msg, + MismatchLine("compressor.ape", "[compress_ratio, head_dim]", cr * hd, cr * 2 * hd)) ) is NOT correct! + values: CHECK( false ) + logged: msg := ... + - attn.compressor.ape: this forward indexes it as [compress_ratio, head_dim] = 2049 elements, the checkpoint carries 4096 + - attn.indexer.wq_b: this forward indexes it as [index_n_heads*index_head_dim, hidden_size] = 2048 elements, the checkpoint carries 1024 +``` + +That second block is the finding itself, printed: the mutated `compressor.ape` +line reads 2049, and the 2048 the OLD bare assertion matched is sitting two lines +below it on `indexer.wq_b`. The old assertion could not have failed. + +**GREEN after**, whole suites rather than the one case: + +| suite | result | +|---|---| +| `test_deepseek_v4_exl3_forward` | 5 cases, 69 assertions, 0 failed | +| `test_deepseek_v4_exl3_loader` | 11 cases, 172 assertions, 0 failed | +| `ctest -R deepseek_v4` | 13/14 passed | + +`test_cuda_deepseek_v4` is the 14th and reports `Not Run`: this build has CUDA +off, so it is not a result either way. + +**Mutations.** Each was verified to have LANDED before the build (a mutation that +never applied reads as a passing test), then built, run, restored, and the restore +verified by SHA-256 before the next one. `ninja rc` and step count are recorded +for each, because a mutation that fails to BUILD silently re-runs the previous +binary and also reads as a pass: + +| # | mutation | ninja | steps | verdict | +|---|---|---|---|---| +| M1 | delete the `comp_wgate` width check | rc=0 | 4/4 | RED (forward) | +| M2 | delete the `comp_ape` width check | rc=0 | 4/4 | RED (forward) | +| M3 | delete the `idx_wq` width check | rc=0 | 4/4 | RED (forward) | +| M4 | delete the `idx_wk` width check | rc=0 | 4/4 | RED (forward) | +| M5 | loader reverts to the collapsed CONSTANT width | rc=0 | 4/4 | RED (both) | +| M6 | **reachability**: delete the production call site | rc=0 | 4/4 | RED (forward) | +| M7 | `RequireDsaDim` accepts ANY width | rc=0 | 4/4 | RED (loader) | +| M8 | loader ALSO accepts the collapsed width | rc=0 | 4/4 | RED (loader) | +| M9 | delete the `comp_norm_weight` check | rc=0 | 4/4 | RED (forward) | +| M10 | delete the `idx_wproj` check | rc=0 | 4/4 | RED (forward) | +| M11 | `indexer.wq_b`'s K also accepts `hidden_size` | rc=0 | 4/4 | RED (loader) | +| M12 | `compressor.ape`'s EXPECTED count wrong on the `cr == 4` layer | rc=0 | 4/4 | **GREEN at `0acf0147f`**, RED after | +| M13 | `compressor.norm.weight`'s EXPECTED count wrong on that layer | rc=0 | 4/4 | **GREEN at `0acf0147f`**, RED after | +| M14 | delete the loader's `coff * index_head_dim` derivation | rc=0 | 4/4 | **GREEN at `0acf0147f`**, RED after | +| — | restored tree | rc=0 | 4/4 | forward 5/69, loader 11/172, `ctest -R deepseek_v4` 13/14 | + +The WHOLE table is re-derived at the branch head, not carried forward from round +1. The round-2 repair changed both the runtime message and the assertions that +read it, and a later commit silently disarming an earlier commit's mutation proof +is a real failure mode — a table measured on a tree that no longer exists proves +nothing about this one. Each mutation was verified to have LANDED before the +build, and each restore verified byte-for-byte by SHA-256 with a rebuild before +the next one, because a mutation that never applied and a mutation that failed to +BUILD both re-run the previous binary and both read as a pass. + +M12 and M13 are scoped to the `cr == 4` layer on purpose. An unscoped version +also breaks the `cr == 128` case in §3, which then reds for a reason that has +nothing to do with the collision — a mutation that reds by collateral damage +measures nothing. The first attempt at M12 did exactly that, and a second attempt +that dropped `cr` from the expression failed to BUILD under `-Werror` +(`ninja rc=1`) while the previous binary still reported SUCCESS, which is the +stale-binary trap this section records step counts for. + +M1-M4 are individually falsifiable only because the refusal reports EVERY mismatch +rather than stopping at the first; a first-mismatch refusal would have made three +of them undetectable. That is also what lets M9 and M10 score at all: deleting one +check leaves the throw but drops that tensor's name from the message. + +**Round 3 (comments and records only).** Focused rebuild `ninja rc=0`, 4/4 steps; +forward 5 cases / 69 assertions, loader 11 cases / 172 assertions. The other ten +`deepseek_v4` suites were RELINKED — `ninja rc=0`, 10/10 steps — rather than run +stale against the previous `libvllm.a`, and `ctest -R deepseek_v4` is 13/14 with +`test_cuda_deepseek_v4` `Not Run` (CUDA off). Every number matches the restored +tree above, which is the whole point: a round that changes no assertion, no +derivation and no call site must not move a count. The two binaries' SHA-256 did +change, and that is expected rather than a build-identity failure — repairing the +carried-half comment added three lines to `deepseek_v4_weights.cpp`, `VT_CHECK` +embeds `__LINE__`, so a comment above a `VT_CHECK` is not a byte-identical build. +The same three lines moved `deepseek_v4_weights.cpp:992` to `:995`, which is the +stale-local-anchor failure arriving once more inside this pull request; every +local anchor in both specs and in the three index rows was re-resolved afterwards. +No mutation was re-run: round 3 touches nothing a mutation scores against. + +## Owed + +- **The full DSA port (option A) has NO owning row.** `MODEL-DSV4-EXL3` carries + it, as its own `## Owed` already says of the dense-MLA policy this supersedes. + It needs the `coff`-overlapped window with `head_offset` role selection + (`common/ops/fused_compress_quant_cache.py:164-183`, the main compressor's + `_fused_kv_compress_norm_rope_insert_sparse_attn`), boundary-only emission, + a compressed KV cache beside a SWA(128) raw cache — which needs the cache + topology [#1960](https://github.com/mudler/vllm.cpp/issues/1960) and + [#1925](https://github.com/mudler/vllm.cpp/issues/1925) are scoping — the + indexer on `qr` over compressed rows, and one joint softmax over the union. + Our `CompressorSaveScoreApe` / `CompressorPoolNorm` primitives are already + generic over width and window and carry over unchanged; the gap is the + composition, not the maths. +- **[#1964](https://github.com/mudler/vllm.cpp/issues/1964) stays open and stays + unfixed here.** The GGUF arm's `dsa_dense` still runs on the real geometry, so + the shipping GGUF DeepSeek-V4 path is still not upstream's attention on 41 of + 43 layers, and the false exactness justification at + `deepseek_v4.cpp:763-775` is still quoted onward by + [#1925](https://github.com/mudler/vllm.cpp/issues/1925). Out of scope by the + dispatch, which excludes changing the GGUF arm's behaviour. +- **`cr == 128` EXL3 layers pass the width check and run the wrong compressor** + (D3). This is [#1976](https://github.com/mudler/vllm.cpp/issues/1976), filed by + the fresh-review repair. It was attributed to #1964 here and that was WRONG: + #1964 is `dsa_dense` routing the GGUF arm to dense MLA, while this is the EXL3 + arm ENTERING the compressor and running a 2-wide pool over the MLA's own latent + (`deepseek_v4.cpp:833`) where upstream pools 128 wide, emits only at boundary + tokens, and reads its own `compressor.wkv` projection. Closing #1964 would not + have closed it, and nothing else tracked it. +- **The `indexer.wq_b` input-space defect.** Upstream projects the indexer query + from `qr`, the q-LoRA latent (`DeepseekV4Indexer.forward`, `attention.py:835`); +our forward feeds it the + hidden state (`deepseek_v4.cpp:915`). After D1 the loader materializes the + tensor at its real `[inh * ihd, q_lora_rank]`, so the forward's `[inh * ihd, H]` + indexing now REFUSES instead of mis-indexing — but the wrong input space is a + real defect at any geometry and option A owns fixing it. +- **No real-checkpoint run.** Every gate here is the hermetic fixture. That the + 99.5 GiB artifact now loads is asserted at the fixture's geometry, not measured + on the artifact, which needs the box and W2 residency. +## Outcome + +**What was measured.** Every upstream shape was confirmed against the checkout at +the pin rather than taken from the scoping spec, asserting UNIQUENESS and not mere +existence: `compressor.py:247-248` (`coff`), `:270-277` (`ape`), `:279-287` +(`fused_wkv_wgate`), `:288` (`RMSNorm(self.head_dim, self.rms_norm_eps)` — the +norm is NOT widened, and it is the file's only `RMSNorm(`), `attention.py:721-726` +and `:835` (`DeepseekV4Indexer.forward`'s `wq_b` on `qr` from `q_lora_rank`), +`:768-776` (the indexer's own +compressor at `index_head_dim`), `:274` (the indexer exists only at `cr == 4`). + +The first cut cited `compressor.py:293` for the norm in five places. `:293` is +`compress_ratio=compress_ratio` inside the `CompressorStateCache` call: the line +drifted by five while the surrounding claim stayed plausible, and it was quoted +onward into the commit body and the append-only index row, where it could not have +been amended after the squash. That is why an anchor is now checked for uniqueness +rather than read once and repeated. + +**It happened a second time, with `attention.py:276`, and the sweep for it stopped +one file short.** Round 1 corrected six copies. A seventh survived in +`deepseek_v4_weights.cpp`'s indexer block, twelve lines above the SAME function's +correct `:274`, and the pull-request body then claimed the correction as complete. +Round 2 corrected the seventh and swept the whole tree, which now carries zero +copies of `:276`. `:274` is `if self.compress_ratio == 4:` and the construct is +unique in `attention.py`; `:276` is a comment about `aux_stream_list`. SIX cited +constructs are NOT unique, and the previous revision of this section said TWO. +The number is MEASURED and not read for: take every upstream anchor this branch +adds (`git diff origin/main...HEAD`, added lines only — all of them land in +`vllm/models/deepseek_v4/`), take the construct quoted beside each, and count that +construct's occurrences in its own file at the pin. + +| construct | file | occurrences | +|---|---|---| +| `self.compressor = DeepseekCompressor(` | `attention.py` | 2 — `:335` (`DeepseekV4Attention.__init__`), `:768` (`DeepseekV4Indexer.__init__`) | +| `self.wq_b(qr)` | `attention.py` | 4 — `:480`, `:514`, `:527` (`DeepseekV4Attention.attention_impl`), `:835` (`DeepseekV4Indexer.forward`) | +| `if (position + 1) % COMPRESS_RATIO != 0:` | `common/ops/fused_compress_quant_cache.py` | 5 — `:164` (`_fused_kv_compress_norm_rope_insert_sparse_attn`), `:364`, `:429`, `:712`, `:891` | +| `head_offset = (tokens >= COMPRESS_RATIO)…* HEAD_SIZE` | `common/ops/fused_compress_quant_cache.py` | 3 — `:182` (that same main-compressor kernel), `:730` (indexer), `:909` (mxfp4 indexer) | +| `swa_only = self.compress_ratio <= 1` | `nvidia/flashinfer_sparse.py` | 3 — `:263` (`DeepseekV4FlashInferMLAAttention.forward_mqa`), `:686`, `:793` | +| `flashinfer_trtllm_batch_decode_sparse_mla_dsv4(` | `nvidia/flashinfer_sparse.py` | 4 — `:486`, `:511`, `:769` (`DeepseekV4FlashInferSM120Attention._forward_decode`), `:888` | + +A seventh is NOT one, and the difference is the whole point of measuring rather +than eyeballing: `[self.coff * self.head_dim, self.coff * self.head_dim],` occurs +at `compressor.py:281` AND `:335`, but the citation is the RANGE `:279-287` and +`:279` is unique, so that anchor already picks out its own line. EVERY anchor in +the table is CORRECT — this is a count defect and not a wrong line — but a +construct that does not pick out its own line cannot be checked by the reader it +was written for, so each is now carried with its enclosing class or kernel named +beside it. + +One more anchor defect of a different kind, found in the same pass: +`save_partial_states.py:85-101` began on a BLANK line. The range now starts at +`:86`, its first real line. + +**Local anchors go stale inside the pull request that writes them, and that is a +separate failure from citing the wrong upstream line.** Five of the six local +`file:line` citations in `dsv4-dsa-geometry.md` were correct at `c00625141`, where +that document was measured, and stale by `0acf0147f`, because #1970's own +implementation moved the lines underneath it. Round 1 swept this document and not +that one. Both are swept now, and every local anchor in both is re-derived at the +branch head. + +**What was rejected, and why.** + +*Accepting two widths.* The first cut derived one width from `coff` and then +relaxed to accept the collapsed one as well, on the premise that refusing it would +delete the synthetic gates. The fresh review disproved the premise: the four +synthetic DSA suites do not load through this arm at all, only the two EXL3 suites +broke, and this change rewrites their fixture regardless. The two-width form was a +divergence from the mirror bought with nothing, and it is gone. What replaced it is +`compress_ratios = {0, 128}` in the forward fixture, where `coff` is 1 and the +collapsed width IS the derived one — the synthetic geometry moves into the RATIO, +where upstream can actually produce it, instead of into the loader's accepted set. + +*Claiming the refusal prevents wrong tokens.* The first cut said `MatVec` had "no +length check" and that a mis-indexed `comp_wgate` was a plausible wrong number. +`deepseek_v4.cpp:413` is an unconditional `VT_CHECK`, both `Gemm` arms check, and +the EXL3 path takes the checked one. The refusal buys a DIAGNOSTIC, and each round +that swept for the claim found copies the round before it had reported gone: six +in round 1, four in round 2, and a twelfth in round 3, in the carried-half comment +of the same file whose corrected paragraph it contradicted. Every copy this branch +can reach — code comment, runtime message, spec, commit body, index row — now says +so; the [#1923](https://github.com/mudler/vllm.cpp/issues/1923) row is on +`origin/main` and append-only, so it keeps the claim permanently and this document +names it instead of claiming a clean tree. + +*Naming every alternative throw in the runtime refusal.* The message says the +alternative is `MatVec weight size mismatch`. That is right for the case the +refusal actually fires on — at the real geometry `comp_wgate`'s `Gemm` runs before +anything reads `comp_ape` or `comp_norm_weight` — but the sentence said "reading +THEM", plural, and a `comp_ape`-only or `comp_norm_weight`-only mismatch throws +`ape size mismatch` / `rms_weight size mismatch` from +`deepseek_v4_compressor.cpp:23,54` instead. The fix is the smallest one that +removes the overclaim: the sentence now names `comp_wgate` and the two other +throws are named in one parenthesis. It does not enumerate every ordering, because +the point is that all three are equally anonymous and that is the whole claim. +Those two line numbers were derived rather than read: `VT_CHECK` embeds +`__LINE__`, and GCC reports the line a multi-line macro invocation BEGINS on, not +the line its message literal sits on — `:24` and `:55` are where the strings are, +`:23` and `:54` are what the throw prints. + +*Two options upstream of this row*, recorded in `dsv4-dsa-geometry.md`: (A) the +full DSA port, correct but a multi-wave model port blocked on cache topology; (B) +a per-layer dense selector, which is still not upstream's attention on 41 layers +and so cannot be gated as parity. + +**Why each default has its value.** `coff = (cr == 4) ? 2 : 1` is upstream's +expression verbatim, not a fit to the artifact. The indexer branch uses the +constant 2 rather than a derivation because the branch is reached only at +`cr == 4`. `indexer.wq_b`'s K is `q_lora_rank` and its refusal cites +`attention.py:721-726` rather than `coff`, because nothing about that tensor is +doubled and the first cut's message blamed the wrong rule. `RequireDsaDim` takes +its `why` per call site for that reason, and all THREE of its derivations now have +a fixture case that reads the message — `coff * head_dim`, +`coff * index_head_dim`, and `q_lora_rank` — because a derivation nothing reads +is a derivation nothing gates. The forward check is +gated on `is_comp || is_indexer` so it fires exactly where a tensor is about to be +read, which is also what leaves the GGUF arm untouched. + +**What this row now owes that it did not before.** +[#1976](https://github.com/mudler/vllm.cpp/issues/1976) exists because D3 was +attributed to #1964 and would have been "closed" by closing an unrelated defect on +another arm. + +## Now + +`ACTIVE` under `MODEL-DSV4-EXL3`. No lifecycle state moved by this document. diff --git a/.agents/specs/model-dsv4-exl3.md b/.agents/specs/model-dsv4-exl3.md index 944cf7f6b..24dbf47cf 100644 --- a/.agents/specs/model-dsv4-exl3.md +++ b/.agents/specs/model-dsv4-exl3.md @@ -888,9 +888,14 @@ compressor width `deepseek_v4.cpp` already documents at the `dsa_dense` comment So the loader REFUSES BY NAME on those three shapes and names the residual. It does not improvise, and it does not widen the host slot to a shape the forward -would then mis-index: `Gemm`'s host arm is a `MatVec` with no length check, so a -`[2*hd, H]` buffer read as `[hd, H]` is a silently wrong number, which is the -`.agents/verification.md` failure this project exists to avoid. +would then mis-index. WITHDRAWN AS WRITTEN by #1970, and recorded here because +this sentence is where the claim started: it said `Gemm`'s host arm is a `MatVec` +with no length check, so a `[2*hd, H]` buffer read as `[hd, H]` is a silently +wrong number. `deepseek_v4.cpp:413` is an unconditional `VT_CHECK` and `Gemm`'s +keep-quant arm checks the shape too, so what that produces is an ANONYMOUS +`vt: MatVec weight size mismatch` naming no tensor and no layer. Refusing by name +buys a DIAGNOSTIC over that, which is still the `.agents/verification.md` concern, +and it is not the difference between wrong tokens and a refusal. **The obvious fix is wrong and the reason is worth recording.** The GGUF arm dodges the same geometry by setting `dsa_dense = (be.gguf != nullptr)` and @@ -1928,6 +1933,30 @@ which is precisely how this landed green locally in the first place. entry is for; it is deliberately NOT attached to [#1923](https://github.com/mudler/vllm.cpp/issues/1923), because that issue is the loader defect and W1c closes it. +- **The real artifact's DSA geometry now LOADS, and the forward REFUSES on it — + [#1970](https://github.com/mudler/vllm.cpp/issues/1970), option C of + [#1961](https://github.com/mudler/vllm.cpp/issues/1961).** This SUPERSEDES the + dense-MLA-policy entry above: the answer is not a shared dense-MLA selector, + because dense MLA is not upstream's attention on a `cr > 0` layer at any + sequence length ([#1964](https://github.com/mudler/vllm.cpp/issues/1964)), so + routing the EXL3 arm there would have been a wrong-but-plausible path rather + than a policy. The loader now derives every DSA width as upstream does + (`coff = 1 + (compress_ratio == 4)`, `vllm/models/deepseek_v4/compressor.py:247-248`) + and `AttentionBlock` refuses BY NAME when a materialized width is not the one + its arithmetic indexes. What stays OWED is the DSA composition itself — the + `coff`-overlapped window with `head_offset` role selection, boundary-only + emission, a compressed KV cache beside a SWA(128) raw cache, the indexer on + `qr` over compressed rows, one joint softmax over the union. **No row owns that + port**; `MODEL-DSV4-EXL3` carries it here until one does, and it needs the + cache topology [#1960](https://github.com/mudler/vllm.cpp/issues/1960) and + [#1925](https://github.com/mudler/vllm.cpp/issues/1925) are scoping. Also owed + and NOT closed by #1970: the GGUF arm's `dsa_dense` still runs the same wrong + attention on 41 of 43 real layers (#1964, excluded from #1970's scope), the + `cr == 128` EXL3 layers pass the width check while their `win = 2` pooling is + still not upstream's 128-wide boundary-emitted compressor, and the + `indexer.wq_b` input-space defect (`x` where upstream uses `qr`) is real at any + geometry. Design, anchors and mutations in + [`specs/dsv4-dsa-loader-accept-forward-refuse.md`](dsv4-dsa-loader-accept-forward-refuse.md). - **Real-checkpoint residency for the coalesced tower — W2.** W1b copies each TP1-coalesced linear into host owner buffers. That is right for the fixture and for W2's byte-parity gate, and it is ~100 GB on the real 216-expert diff --git a/include/vllm/model_executor/models/deepseek_v4.h b/include/vllm/model_executor/models/deepseek_v4.h index 4a655ab00..46b3cbb02 100644 --- a/include/vllm/model_executor/models/deepseek_v4.h +++ b/include/vllm/model_executor/models/deepseek_v4.h @@ -159,14 +159,23 @@ struct DeepseekV4LayerHostWeights { std::vector attn_sink; // [n_heads] std::vector wo_a; // [n_groups, o_lora_rank, in_per_group] std::vector wo_b; // [H, n_groups*o_lora_rank] - // DSA Lightning-Indexer (indexer layers only; empty otherwise). - std::vector idx_wq; // [index_n_heads*index_head_dim, H] - std::vector idx_wk; // [index_head_dim, H] - std::vector idx_wproj; // [index_n_heads, H] - // DSA compressor (compressor layers only; empty otherwise). - std::vector comp_wgate; // [head_dim, H] (produces the pool score) - std::vector comp_ape; // [compress_ratio, head_dim] - std::vector comp_norm_weight; // [head_dim] + // DSA compressor + Lightning-Indexer (those layers only; empty otherwise). + // + // TWO GEOMETRIES MEET IN THESE SLOTS, and the shapes below are the LOADED ones + // (#1970). The EXL3 loader materializes each at the width upstream DERIVES for + // the layer — `coff = 1 + (compress_ratio == 4)`, `compressor.py:247-248`, and + // `wq_b`'s natural `q_lora_rank` K, `attention.py:721-726`. `AttentionBlock` + // indexes the COLLAPSED synthetic geometry instead (`comp_wgate` as + // `[head_dim, H]`, `comp_ape` as `[compress_ratio, head_dim]`, `idx_wq` as + // `[index_n_heads*index_head_dim, H]`, `idx_wk` as `[index_head_dim, H]`), so + // where the two differ it REFUSES BY NAME rather than reading either. They + // coincide exactly where `coff` is 1 — every `compress_ratio != 4` layer. + std::vector idx_wq; // [index_n_heads*index_head_dim, q_lora_rank] + std::vector idx_wk; // [coff*index_head_dim, H] + std::vector idx_wproj; // [index_n_heads, H] (not widened upstream) + std::vector comp_wgate; // [coff*head_dim, H] (the pool score) + std::vector comp_ape; // [compress_ratio, coff*head_dim] + std::vector comp_norm_weight; // [head_dim] (compressor.py:288) // MoE router: learned gate + (non-hash) noaux_tc bias OR (hash) tid2eid table. std::vector gate_weight; // [n_routed_experts, H] std::vector gate_bias; // [n_routed_experts] (non-hash layers) diff --git a/src/vllm/model_executor/models/deepseek_v4.cpp b/src/vllm/model_executor/models/deepseek_v4.cpp index 0a0c55970..59d35eaad 100644 --- a/src/vllm/model_executor/models/deepseek_v4.cpp +++ b/src/vllm/model_executor/models/deepseek_v4.cpp @@ -645,6 +645,105 @@ std::vector Slice(const std::vector& v, int64_t off, int64_t len) } // ── 512-wide MLA attention block (W3 + W4 primitives) : [T,H] -> [T,H] ──────── +// #1970 — THE LENGTH CHECK THE DSA PATH NEVER HAD. +// +// `AttentionBlock` indexes the DSA tensors at the COLLAPSED synthetic geometry: +// `comp_wgate` as [head_dim, hidden_size], `comp_ape` as [compress_ratio, +// head_dim], `idx_wq` as [index_n_heads*index_head_dim, hidden_size], `idx_wk` +// as [index_head_dim, hidden_size]. Since #1970 the EXL3 loader materializes +// them at the width the REAL artifact stores — upstream's +// `coff = 1 + (compress_ratio == 4)` (vllm/models/deepseek_v4/compressor.py:247-248) +// — so the two can now disagree. +// +// AND A DISAGREEMENT HERE IS ANONYMOUS, NOT SILENT. Be exact about what this +// buys, because overstating it is the defect #1964 was filed for. `Gemm`'s host +// arm is a `MatVec` whose size assertion is UNCONDITIONAL — `deepseek_v4.cpp:413` +// is a plain `VT_CHECK`, a throw rather than an `assert`, so `NDEBUG` does not +// remove it — and its keep-quant arm checks the shape too. A [2*head_dim, +// hidden_size] weight read at a [head_dim, hidden_size] stride therefore does NOT +// produce a plausible wrong number. It throws +// +// vt: MatVec weight size mismatch at deepseek_v4.cpp:413 +// +// which names no tensor, no layer, no geometry and no missing capability, from +// the middle of a forward, on a checkpoint that loaded successfully. +// +// So this is a DIAGNOSTICS improvement, and that is the whole of it: it replaces +// an anonymous crash with a precise named refusal, listing EVERY mismatched +// tensor with both counts and naming the composition that is missing. It is not +// the difference between wrong tokens and a refusal, and it must not be described +// as one. +// +// EVERY mismatch is collected and reported together, not just the first. A +// refusal that stopped at the first would make the remaining checks +// unfalsifiable: deleting any one of them would still throw on an earlier one, +// so a mutation could not tell a live check from a dead one. +void RequireDsaGeometryOrRefuse(const DeepseekV4LayerHostWeights& L, + const DeepseekV4Params& p, int64_t layer, + bool is_comp, bool is_indexer) { + const int64_t H = p.hidden_size; + const int64_t hd = p.head_dim; + std::string bad; + auto want = [&](const char* tensor, const char* indexed_as, size_t got, + int64_t expect) { + if (static_cast(got) == expect) return; + bad += "\n - attn." + std::string(tensor) + ": this forward indexes it as " + + indexed_as + " = " + std::to_string(expect) + + " elements, the checkpoint carries " + std::to_string(got); + }; + if (is_comp) { + const int64_t cr = p.compress_ratio(layer); + want("compressor.ape", "[compress_ratio, head_dim]", L.comp_ape.size(), cr * hd); + want("compressor.wgate.weight", "[head_dim, hidden_size]", L.comp_wgate.size(), + hd * H); + want("compressor.norm.weight", "[head_dim]", L.comp_norm_weight.size(), hd); + } + if (is_indexer) { + const int64_t inh = p.index_n_heads; + const int64_t ihd = p.index_head_dim; + want("indexer.wq_b", "[index_n_heads*index_head_dim, hidden_size]", + L.idx_wq.size(), inh * ihd * H); + want("indexer.compressor.wkv.weight", "[index_head_dim, hidden_size]", + L.idx_wk.size(), ihd * H); + want("indexer.weights_proj.weight", "[index_n_heads, hidden_size]", + L.idx_wproj.size(), inh * H); + } + VT_CHECK( + bad.empty(), + std::string("DeepseekV4 forward: REFUSING the DSA path on layer ") + + std::to_string(layer) + + " — the checkpoint carries this layer's DSA tensors at a geometry this " + "forward does not implement. Reading the widened `comp_wgate` at the " + "width it DOES index throws an anonymous `MatVec weight size mismatch` " + "from inside the forward (deepseek_v4.cpp:413) that names none of this. " + "(That is the message the REAL geometry produces, because `comp_wgate`'s " + "Gemm runs first. A `comp_ape`- or `comp_norm_weight`-only mismatch " + "instead throws `ape size mismatch` / `rms_weight size mismatch` from " + "CompressorSaveScoreApe / CompressorPoolNorm " + "(deepseek_v4_compressor.cpp:23,54) — equally anonymous.) Refusing on:" + bad + + "\n WHAT IS MISSING: upstream's DSA composition. The extra width is " + "`coff = 1 + (compress_ratio == 4)` " + "(vllm/models/deepseek_v4/compressor.py:247-248), and its two halves " + "are the two OVERLAPPING compression windows a token belongs to — a " + "role a row acquires only relative to the window gathering it " + "(common/ops/fused_compress_quant_cache.py:164-183), never recoverable " + "from the tensor alone. Reaching it needs the coff-overlapped window " + "with head_offset role selection, emission at boundary tokens only " + "((position + 1) % compress_ratio == 0) into a SEPARATE compressed KV " + "cache beside a sliding-window raw cache, and the indexer's query " + "projected from `qr` (q_lora_rank) instead of the hidden state " + "(vllm/models/deepseek_v4/attention.py:721-726, :835). None of that is " + "implemented here, and dense MLA is NOT a substitute for it at any " + "sequence length " + "(https://github.com/mudler/vllm.cpp/issues/1964).\n" + " The loader accepts this geometry ON PURPOSE, so every NON-DSA " + "capability of the artifact is reachable rather than blocked behind a " + "path none of them use (MODEL-DSV4-EXL3 option C, " + "https://github.com/mudler/vllm.cpp/issues/1970). The DSA port itself " + "is OWED and has no owning row — see " + ".agents/specs/dsv4-dsa-loader-accept-forward-refuse.md `## Owed`."); +} + std::vector AttentionBlock(const DeepseekV4LayerHostWeights& L, const DeepseekV4GgufLayerWeights* Lq, const DeepseekV4Params& p, @@ -678,6 +777,13 @@ std::vector AttentionBlock(const DeepseekV4LayerHostWeights& L, const bool is_indexer = p.has_indexer(layer) && !dsa_dense; const bool is_comp = p.has_compressor(layer) && !dsa_dense; + // Gated on the predicates above, so it fires only where this forward is about + // to READ one of these tensors. A layer that does not enter the DSA path reads + // none of them, and the GGUF arm (`dsa_dense`) enters it on no layer at all — + // so that arm's behaviour is byte for byte unchanged by #1970, and its own + // separate defect stays owed under #1964. + RequireDsaGeometryOrRefuse(L, p, layer, is_comp, is_indexer); + // 1. q [T,nh,hd] and raw kv latent [T,hd] (num_key_value_heads=1 MLA). The MLA // linears (wq_a, wq_b, wkv) run the keep-quant GEMM (Gemm) — the whole batch // at once — then the per-token RMSNorm(q_norm/kv_norm) + per-head RoPE. diff --git a/src/vllm/model_executor/models/deepseek_v4_weights.cpp b/src/vllm/model_executor/models/deepseek_v4_weights.cpp index db94fee35..0840a4155 100644 --- a/src/vllm/model_executor/models/deepseek_v4_weights.cpp +++ b/src/vllm/model_executor/models/deepseek_v4_weights.cpp @@ -336,9 +336,15 @@ Exl3RankSlice ReadRankSlice(const StIndex& index, const std::string& base, int b // that tower. // // Each reader is given the EXPECTED shape, derived from the resolved config, and -// refuses a mismatch BY NAME. That is not defensive decoration: `Gemm`'s host arm -// is a `MatVec` with no length check, so a tensor materialized at the wrong shape -// is a silently wrong number rather than a crash. +// refuses a mismatch BY NAME. That is not defensive decoration, and the reason is +// DIAGNOSTIC rather than numeric. A tensor materialized at the wrong shape does +// not produce a wrong number: `Gemm`'s host arm is a `MatVec` whose size +// assertion is unconditional (`deepseek_v4.cpp:413`, a plain `VT_CHECK` and not +// an `assert`, so it survives `NDEBUG`), and its keep-quant arm checks too. What +// it produces is an ANONYMOUS throw — `vt: MatVec weight size mismatch at +// deepseek_v4.cpp:413` — that names neither the tensor, nor the layer, nor the +// geometry, nor what is missing. Refusing HERE replaces that with a message the +// reader can act on. // The carried half's own quantization recipe. The artifact records it twice — // `quantization_config.base_quantization_config` and the top-level @@ -404,6 +410,51 @@ void RequireShape(const StTensor& t, const std::vector& want, "mis-index (" + kExl3Row + " W1c)"); } +// #1970 — the DSA family's width is DERIVED the way upstream derives it, and +// the checkpoint is required to AGREE. +// +// `coff = 1 + (compress_ratio == 4)` (`vllm/models/deepseek_v4/compressor.py:247-248`) +// is a pure function of `compress_ratio`. It sizes the APE table (`:270-277`), +// BOTH halves of the fused `wkv|wgate` projection (`:279-287`) and the compressed +// state cache (`:291`); the norm is NOT widened (`:288` is +// `RMSNorm(self.head_dim, self.rms_norm_eps)`). So for a given `compress_ratio` +// upstream can emit exactly ONE width, and a `cr == 4` checkpoint carrying an +// UNDOUBLED family is a checkpoint upstream cannot load at all. +// +// This therefore derives the width and refuses ANY other BY NAME. It does not +// also accept a "collapsed" width: `AGENTS.md` requires this loader to mirror +// every mode upstream defines, so accepting a width upstream can never emit is a +// divergence, and production acceptance must not be widened to suit a fixture. +// A synthetic geometry belongs in the fixture's `compress_ratio` — at `cr == 128` +// `coff` is 1 and the derived width IS the collapsed one — not in the set of +// shapes the loader will take from a real artifact. +// +// `why` names the derivation in the refusal, because the three call sites derive +// their width from three DIFFERENT rules and a message that named only `coff` +// would be wrong on the one that has nothing to do with it. +// +// EACH of the three needs its OWN fixture case, and two of them did not get one +// until a mutation went looking. This check is strictly weaker than the +// `RequireShape` that follows it, so its only product is the derivation in `why` +// — and a `why` no fixture ever reads is not gated. A wholly collapsed +// checkpoint refuses on the FIRST derivation and never reaches the other two, so +// `collapsed_indexer_wq_b` and `collapsed_indexer_wkv` exist to write the real +// geometry everywhere else and collapse exactly one tensor. +void RequireDsaDim(const std::vector& shape, size_t dim, int64_t want, + const std::string& name, const std::string& why) { + VT_CHECK(dim < shape.size(), + std::string("deepseek-v4 exl3 loader: ") + name + " must have at least " + + std::to_string(dim + 1) + " dimensions, got " + ShapeText(shape) + + " (" + kExl3Row + " W1c)"); + VT_CHECK(shape[dim] == want, + std::string("deepseek-v4 exl3 loader: ") + name + " dimension " + + std::to_string(dim) + " must be " + std::to_string(want) + " — " + why + + " — got " + std::to_string(shape[dim]) + " in " + ShapeText(shape) + + ". This is the width upstream DERIVES for this layer, and it is the " + "only one upstream emits; refusing rather than materializing a family " + "the forward would mis-index (" + kExl3Row + " W1c)"); +} + // The MATERIALIZING counterpart of W1b's `require`. Everything it reads is // accounted exactly as before; the difference is that the bytes now land in the // host tower. @@ -524,6 +575,13 @@ class Exl3CarriedReader { // hunt for the missing slot. void Account(const std::string& name) { (void)Take(name); } + // The stored shape, WITHOUT accounting the tensor. The DSA family's width has + // to be read before the family can be required to agree with itself, and the + // tensor is then taken normally by `Float`/`Fp8Block` below. + const std::vector& PeekShape(const std::string& name) { + return RequireTensor(index_, name)->shape; + } + private: static int64_t Numel(const std::vector& s) { int64_t n = 1; @@ -789,8 +847,11 @@ DeepseekV4Weights LoadDeepseekV4Exl3(const std::vector& shards, // routed-expert block EXL3 replaced — MATERIALIZED into the host-float // tower `ForwardComposeImpl` composes with (W1c). W1b only counted these. // Every destination shape comes from the resolved config, so the refusal - // that fires on the real artifact's DSA geometry names the tensor rather - // than producing a wrong number (see `## W1c design` W1c-4). ──────────── + // that fires on the real artifact's DSA geometry NAMES the tensor instead + // of the ANONYMOUS `vt: MatVec weight size mismatch` a wrong shape throws + // anyway. That is a DIAGNOSTIC and the whole of it, not the difference + // between wrong tokens and a refusal — the reader-shape block above says + // why (see `## W1c design` W1c-4). ───────────────────────────────────── std::unordered_set routed; const Exl3CarriedFp8Recipe recipe = ResolveCarriedFp8Recipe(config); Exl3CarriedReader carried(index, recipe, &routed, &accounted); @@ -853,22 +914,85 @@ DeepseekV4Weights LoadDeepseekV4Exl3(const std::vector& shards, if (p.has_compressor(l)) { const int64_t cr = p.compress_ratio(l); - hl.comp_ape = carried.Float(a + "compressor.ape", {cr, hd}); + // UPSTREAM'S OWN EXPRESSION, and the only thing the doubled width comes + // from (#1970, scoped by #1961): + // + // vllm/models/deepseek_v4/compressor.py:247-248 + // self.overlap = compress_ratio == 4 + // self.coff = 1 + self.overlap + // + // The width is DERIVED, not chosen from a set: every tensor of the family + // is required at `coff * head_dim`, so a half-widened checkpoint refuses on + // whichever member disagrees. At `cr == 128` `overlap` is false, `coff` is + // 1, and the derived width is the collapsed one — which is why 20 of the + // real artifact's 41 compressor layers already loaded before #1970. + // + // MATERIALIZING THESE IS NOT THE SAME AS BEING ABLE TO USE THEM. The two + // halves are the two OVERLAPPING compression windows a token belongs to, + // and which half a row plays is decided at gather time by window position + // (`common/ops/fused_compress_quant_cache.py:164-183`), never recoverable + // from the tensor alone. The loader accepts them so every NON-DSA + // capability of the artifact becomes reachable; `AttentionBlock` refuses + // BY NAME rather than indexing them at a width they do not have. + const int64_t coff = (cr == 4) ? 2 : 1; + const int64_t cw = coff * hd; + const std::string wg = a + "compressor.wgate.weight"; + RequireDsaDim(carried.PeekShape(wg), 0, cw, wg, + "`coff * head_dim`, where `coff = 1 + (compress_ratio == 4)` " + "(compressor.py:247-248) and this layer's compress_ratio is " + + std::to_string(cr)); + // `norm.weight` is NOT widened by `coff`: upstream is + // `RMSNorm(self.head_dim, self.rms_norm_eps)` (`compressor.py:288`). + hl.comp_ape = carried.Float(a + "compressor.ape", {cr, cw}); hl.comp_norm_weight = carried.Float(a + "compressor.norm.weight", {hd}); - hl.comp_wgate = carried.Float(a + "compressor.wgate.weight", {hd, H}); - // Accounted, no destination: the compressor's KV is the MLA's own `kraw` - // latent in `AttentionBlock`, so no host slot reads a separate projection. + hl.comp_wgate = carried.Float(wg, {cw, H}); + // Accounted, no destination: the collapsed-geometry compressor reuses the + // MLA's own `kraw` latent as its KV, so no host slot reads a separate + // projection. Upstream HAS one, and wiring it is part of the owed DSA + // composition rather than of this wave. carried.Account(a + "compressor.wkv.weight"); } if (p.has_indexer(l)) { - // The indexer's own compressor is accounted whole: `L.idx_wk` is its KV - // projection and the other three have no host destination at this - // geometry, the same way the main compressor's KV has none. - hl.idx_wk = carried.Float(a + "indexer.compressor.wkv.weight", {ihd, H}); + // The indexer exists ONLY at `cr == 4` (`attention.py:274`), so upstream's + // `coff` here is always 2. It carries its OWN `DeepseekCompressor` at + // `head_dim = index_head_dim` with the same ratio (`attention.py:768-776`), + // so the same rule widens its family. + // + // `L.idx_wk` is its KV projection; the other three have no host + // destination at this geometry, the same way the main compressor's KV has + // none. + const std::string ik = a + "indexer.compressor.wkv.weight"; + const int64_t iw = 2 * ihd; + RequireDsaDim(carried.PeekShape(ik), 0, iw, ik, + "`coff * index_head_dim`, and the indexer exists only at " + "`compress_ratio == 4` (attention.py:274) where `coff` is 2 " + "(compressor.py:247-248)"); + hl.idx_wk = carried.Float(ik, {iw, H}); for (const char* c : {"ape", "norm.weight", "wgate.weight"}) carried.Account(a + "indexer.compressor." + c); hl.idx_wproj = carried.Float(a + "indexer.weights_proj.weight", {inh, H}); - hl.idx_wq = carried.Fp8Block(a + "indexer.wq_b", inh * ihd, H); + // NOT A WIDTH PROBLEM AT ALL, and worth separating from the rest. Upstream + // builds this as `ReplicatedLinear(self.q_lora_rank, self.head_dim * + // self.n_head)` (`attention.py:721-726`) and calls it on `qr`, the q-LoRA + // latent (`:835`). Its K is therefore `q_lora_rank`, and the stored + // `[inh*ihd, q_lora_rank]` is at its NATURAL size with nothing doubled. + // Our forward feeds it the HIDDEN STATE, which is the wrong input space at + // ANY geometry. No gate ever saw it because the collapsed fixture WRITES + // `wq_b` at `K = H` to match our forward — `H` and `q_lora_rank` do not + // coincide there (`dsv4_exl3_fixture.h`: `kHidden` is 256, `kQLora` is + // 128). Owed — see + // `.agents/specs/dsv4-dsa-loader-accept-forward-refuse.md` `## Owed`. + // + // Its K is required at `q_lora_rank` and the refusal says so. It is NOT a + // `coff` width and citing one here would be wrong: nothing about this + // tensor is doubled. + const std::string iq = a + "indexer.wq_b"; + RequireDsaDim(carried.PeekShape(iq + ".weight"), 1, qlr, iq, + "`q_lora_rank`, because upstream builds `wq_b` as " + "`ReplicatedLinear(q_lora_rank, head_dim * n_head)` " + "(attention.py:721-726) and calls it on `qr` (:835); this is " + "the tensor's NATURAL size and no `coff` applies to it"); + hl.idx_wq = carried.Fp8Block(iq, inh * ihd, qlr); } const std::string f = b + "ffn."; diff --git a/tests/vllm/models/dsv4_exl3_fixture.h b/tests/vllm/models/dsv4_exl3_fixture.h index c4d388b92..def01ce56 100644 --- a/tests/vllm/models/dsv4_exl3_fixture.h +++ b/tests/vllm/models/dsv4_exl3_fixture.h @@ -310,9 +310,55 @@ struct FixtureOptions { // Drop `base_quantization_config`, so the loader has no recipe for the // carried FP8 half and must refuse rather than assume a block size. bool omit_base_quant_config = false; - // Write the compressor at the REAL artifact's width (`2 * head_dim`, - // ds4 `coff = 2`) instead of the collapsed one the host forward indexes. - bool real_compressor_width = false; + // Write the DSA family at the REAL artifact's PER-LAYER geometry instead of + // the collapsed one the host forward indexes (#1970). + // + // Upstream derives the whole thing from one line, + // `coff = 1 + (compress_ratio == 4)` + // (`vllm/models/deepseek_v4/compressor.py:247-248` at the parity pin + // `5559679229bc961848b121ccdeaa8fa5d79bec98`), and spends it on the APE table + // (`:270-277`) and on the fused `wkv|wgate` projection (`:279-287`) — and NOT + // on the norm, which is `RMSNorm(self.head_dim, self.rms_norm_eps)` (`:288`). + // The indexer carries + // its OWN compressor at `head_dim = index_head_dim` and the same ratio + // (`attention.py:768-776`), so its family doubles too; and its `wq_b` is + // `ReplicatedLinear(q_lora_rank, head_dim * n_head)` (`:721-726`), whose K is + // `q_lora_rank` rather than `hidden_size` — not a width at all, but the + // q-LoRA input space our forward does not project from. + // + // This REPLACES `real_compressor_width`, which doubled UNCONDITIONALLY. The + // one case that used it wrote a doubled compressor on a `cr == 128` layer, + // where upstream's `coff` is 1 — a width the artifact does not store and the + // pin does not produce. Keying on `compress_ratio` is what stops this fixture + // from describing a geometry no oracle would emit. + // + // LEAVING THIS FALSE AT `cr == 4` IS NOT A NEUTRAL DEFAULT. It writes the + // collapsed, UNDOUBLED family, which upstream's `coff` cannot produce at that + // ratio and which the loader therefore refuses (#1970). A case that wants a + // compressor layer the host forward can actually run uses `cr == 128`, where + // `coff` is 1 and the derived width IS the collapsed one. + bool real_dsa_geometry = false; + // A THIRD width, which is neither upstream's `coff` width nor the collapsed + // one. The loader DERIVES one width and refuses everything else, and a flag + // that only ever toggled between `coff` and collapsed could not tell a derived + // rule from a two-value allow-list — so this writes a width no oracle produces + // and no forward indexes, and the loader must still refuse it BY NAME. + bool bogus_dsa_width = false; + // `indexer.wq_b` at `K = hidden_size` while the REST of the family is at the + // real geometry. Its K is `q_lora_rank` upstream + // (`attention.py:721-726`, called on `qr` at `:835`) and is NOT a `coff` + // width, so it is derived by a DIFFERENT rule from every other DSA tensor and + // needs its own case: without this the compressor refuses first and the + // `wq_b` check is never reached, which makes it unfalsifiable. + // `kHidden` (256) and `kQLora` (128) differ, so this is a real distinction. + bool collapsed_indexer_wq_b = false; + // The indexer's OWN `DeepseekCompressor` KV projection at the COLLAPSED + // `index_head_dim` while the rest of the family is at the real geometry. + // Without it the main compressor refuses FIRST and the loader's + // `coff * index_head_dim` derivation is never read, which leaves that + // derivation's message unfalsifiable — deleting the check kept both suites + // green. Same shape of hole as `collapsed_indexer_wq_b`, one tensor over. + bool collapsed_indexer_wkv = false; int64_t compress_ratio(int layer) const { return layer < static_cast(compress_ratios.size()) @@ -321,6 +367,12 @@ struct FixtureOptions { } bool has_compressor(int layer) const { return compress_ratio(layer) != 0; } bool has_indexer(int layer) const { return compress_ratio(layer) == 4; } + // `compressor.py:247-248`, verbatim. 1 unless the real geometry is requested + // AND this layer is one of the overlapping `cr == 4` ones. + int64_t coff(int layer) const { + if (bogus_dsa_width) return 3; // neither width; must refuse + return (real_dsa_geometry && compress_ratio(layer) == 4) ? 2 : 1; + } bool is_hash_layer(int layer) const { return layer < num_hash_layers; } }; @@ -460,25 +512,44 @@ inline std::vector CarriedEntries(const FixtureOptions& opt) { push(F32Entry(a + "attn_sink", {kHeads}, 0.7f, 0.0f)); if (opt.has_compressor(l)) { - // The REAL artifact stores the compressor at `2 * head_dim` - // (ds4 `coff = 2`); the collapsed geometry the host forward indexes is - // `head_dim`. `real_compressor_width` writes the former on purpose, to - // gate the loader's refusal. - const int64_t cw = opt.real_compressor_width ? 2 * kHeadDim : kHeadDim; + // The REAL artifact stores the compressor family at `coff * head_dim`, + // where `coff = 1 + (compress_ratio == 4)` + // (`vllm/models/deepseek_v4/compressor.py:247-248`); the collapsed + // synthetic geometry the host forward indexes is `head_dim`. + // `real_dsa_geometry` writes the former ON THE cr == 4 LAYERS ONLY, which + // is where upstream's `overlap` is true and where the real artifact's four + // refusing tensors live. + // + // `norm.weight` is DELIBERATELY not widened: upstream is + // `RMSNorm(self.head_dim, self.rms_norm_eps)` (`:288`), so a doubled norm + // would be a shape the artifact does not carry and the loader must not + // learn to expect. + const int64_t cw = opt.coff(l) * kHeadDim; push(F32Entry(a + "compressor.ape", {opt.compress_ratio(l), cw}, 0.2f, 0.0f)); - push(Bf16Entry(a + "compressor.norm.weight", {cw}, 0.1f, 1.0f)); + push(Bf16Entry(a + "compressor.norm.weight", {kHeadDim}, 0.1f, 1.0f)); push(Bf16Entry(a + "compressor.wgate.weight", {cw, H}, 0.3f, 0.0f)); push(Bf16Entry(a + "compressor.wkv.weight", {cw, H}, 0.3f, 0.0f)); } if (opt.has_indexer(l)) { const int64_t ihd = opt.index_head_dim; const int64_t inh = opt.index_n_heads; - push(F32Entry(a + "indexer.compressor.ape", {4, ihd}, 0.2f, 0.0f)); + // The indexer's own `DeepseekCompressor` runs at `head_dim = + // index_head_dim` with the SAME ratio (`attention.py:768-776`), so the + // same `coff` applies to its family and its norm is likewise undoubled. + const int64_t iw = opt.coff(l) * ihd; + push(F32Entry(a + "indexer.compressor.ape", {4, iw}, 0.2f, 0.0f)); push(Bf16Entry(a + "indexer.compressor.norm.weight", {ihd}, 0.1f, 1.0f)); - push(Bf16Entry(a + "indexer.compressor.wgate.weight", {ihd, H}, 0.3f, 0.0f)); - push(Bf16Entry(a + "indexer.compressor.wkv.weight", {ihd, H}, 0.3f, 0.0f)); + push(Bf16Entry(a + "indexer.compressor.wgate.weight", {iw, H}, 0.3f, 0.0f)); + push(Bf16Entry(a + "indexer.compressor.wkv.weight", + {opt.collapsed_indexer_wkv ? ihd : iw, H}, 0.3f, 0.0f)); push(Bf16Entry(a + "indexer.weights_proj.weight", {inh, H}, 0.3f, 0.0f)); - push_all(Fp8BlockEntries(a + "indexer.wq_b", inh * ihd, H)); + // NOT a width. `wq_b` is `ReplicatedLinear(q_lora_rank, head_dim * + // n_head)` (`attention.py:721-726`) called on `qr` (`:835`), so its K is + // `q_lora_rank`. The collapsed fixture writes `H` because our forward + // feeds it the hidden state. + push_all(Fp8BlockEntries( + a + "indexer.wq_b", inh * ihd, + (opt.real_dsa_geometry && !opt.collapsed_indexer_wq_b) ? kQLora : H)); } const std::string f = b + "ffn."; diff --git a/tests/vllm/models/test_deepseek_v4_exl3_forward.cpp b/tests/vllm/models/test_deepseek_v4_exl3_forward.cpp index 627a0ff08..2c86b1331 100644 --- a/tests/vllm/models/test_deepseek_v4_exl3_forward.cpp +++ b/tests/vllm/models/test_deepseek_v4_exl3_forward.cpp @@ -68,21 +68,48 @@ namespace { // The shape of the model the forward fixture describes: two layers so the MoE // runs more than once, layer 0 hash-routed and layer 1 carrying the DSA -// compressor + Lightning-Indexer, so the load has to materialize every carried -// family the host forward reads. `topk = 2` over the fixture's two routed -// experts keeps both live on every token. +// compressor, so the load has to materialize every carried family the host +// forward reads. `topk = 2` over the fixture's two routed experts keeps both +// live on every token. +// +// WHY LAYER 1 IS `cr == 128` AND NOT `cr == 4` (#1970). The loader derives the +// compressor width from `coff = 1 + (compress_ratio == 4)` +// (`vllm/models/deepseek_v4/compressor.py:247-248`), which at `cr == 4` is 2 — +// so a `cr == 4` layer the host forward can RUN would have to be written at the +// collapsed, undoubled width, and that is a checkpoint upstream cannot emit and +// this loader refuses. At `cr == 128` `coff` is 1, the derived width and the +// collapsed one are the same value, and these cases keep driving an EXL3-loaded +// compressor layer end to end through the production forward. +// +// WHAT THAT COSTS, stated rather than implied: no case here runs an EXL3-loaded +// INDEXER forward, because the indexer exists only at `cr == 4` +// (`attention.py:274`) where the real artifact's geometry is the one the forward +// refuses. The indexer maths itself stays gated at the synthetic geometry by +// `test_deepseek_v4_forward.cpp` and `test_deepseek_v4_dsa.cpp`, which do not +// load through this arm. FixtureOptions ForwardFixtureOptions() { FixtureOptions opt; opt.layers = 2; opt.num_hash_layers = 1; opt.topk = 2; - opt.compress_ratios = {0, 4}; + opt.compress_ratios = {0, 128}; opt.index_n_heads = 2; opt.index_head_dim = 4; opt.index_topk = 3; return opt; } +// The same model with layer 1 at `cr == 4` and the DSA family written at the +// width the REAL DeepSeek-V4-Flash artifact stores: `coff` is 2 there, so the +// compressor + indexer families are doubled and `indexer.wq_b`'s K is +// `q_lora_rank`. This LOADS and the forward REFUSES it. +FixtureOptions RealDsaFixtureOptions() { + FixtureOptions opt = ForwardFixtureOptions(); + opt.compress_ratios = {0, 4}; + opt.real_dsa_geometry = true; + return opt; +} + struct Rng { uint32_t s = 0x243F6A88u; float next(float scale) { @@ -173,6 +200,25 @@ struct QueueGuard { QueueGuard& operator=(const QueueGuard&) = delete; }; +// One line of `RequireDsaGeometryOrRefuse`'s mismatch list, rendered exactly as +// the refusal renders it (`deepseek_v4.cpp`, the `want` lambda). +// +// WHY THE COUNTS ARE ASSERTED BOUND TO A TENSOR NAME AND NOT ON THEIR OWN. At +// this fixture's dimensions several of the counts COINCIDE: `compress_ratio * +// head_dim`, `index_n_heads * index_head_dim * hidden_size` and `2 * +// index_head_dim * hidden_size` are all 2048, and `head_dim - 1` and +// `index_n_heads * hidden_size - 1` are both 511. A bare `Mentions(msg, +// "2048")` therefore does not say which tensor reported it, and the round-2 +// fresh review measured that directly: a mutation that made `compressor.ape`'s +// expected count wrong left both suites GREEN, because another tensor's line +// still carried the number. Binding the count to the name closes that. +std::string MismatchLine(const char* tensor, const char* indexed_as, int64_t indexed, + int64_t carried) { + return "attn." + std::string(tensor) + ": this forward indexes it as " + indexed_as + + " = " + std::to_string(indexed) + " elements, the checkpoint carries " + + std::to_string(carried); +} + const std::vector kTokens = {3, 7, 1}; const std::vector kPositions = {0, 1, 2}; @@ -347,3 +393,166 @@ TEST_CASE("dsv4 exl3 W2d: the FUSED MoE op is what the LOADED forward dispatches " per_expert_calls=", per_expert, " vs dequantized-dense rel_rms=", equiv); CHECK(equiv <= 2.0e-2); } + +TEST_CASE("dsv4 exl3 #1970: the REAL DSA geometry LOADS and the FORWARD refuses by name") { + // OPTION C of `.agents/specs/dsv4-dsa-geometry.md` (#1961), gated end to end: + // the loader materializes every DSA tensor at the width the REAL + // DeepSeek-V4-Flash artifact stores, and `AttentionBlock` refuses BY NAME + // instead of indexing one at a width it does not have. + // + // WHY IT ENTERS THROUGH THE LOADER AND THE MODEL, AND NOT THROUGH THE TYPE. + // `.agents/reachability.md`: a test that constructs `DeepseekV4Weights` by + // hand proves the check works and NEVER that anything reaches it. This row has + // already paid for that once — W2's reachability claim was gated on a struct + // no loader could produce, and the real `vllm-server` load then generated zero + // tokens (#1923). So this drives `vllm::LoadDeepseekV4ForCausalLMWeights`, the + // entry `deepseek_v4_registry.cpp` routes `ModelRegistry::Load` to, and + // `vllm::DeepseekV4Model::Forward`, what `ForwardDeepseekV4ForCausalLM` calls. + auto f = BuildFixture(RealDsaFixtureOptions()); + + // 1. IT LOADS. Before #1970 `RequireShape` refused four tensors here, and on + // the real artifact that is 41 of 43 layers — so the EXL3 tower at real + // scale, MoE, MTP and W2 residency were all unreachable behind a geometry + // none of them use. + const DeepseekV4Weights w = + vllm::LoadDeepseekV4ForCausalLMWeights(f->shards, f->config); + REQUIRE(w.has_exl3_weights); + REQUIRE(w.has_host_weights); + + // ...and it materialized the REAL widths, not the collapsed ones. A loader + // that "accepted" by silently truncating to `head_dim` would pass the refusal + // check below for the wrong reason, so the widths are asserted directly. + const int64_t hd = w.params.head_dim; + const int64_t H = w.params.hidden_size; + const int64_t cr = w.params.compress_ratio(1); + const int64_t inh = w.params.index_n_heads; + const int64_t ihd = w.params.index_head_dim; + REQUIRE(cr == 4); + const DeepseekV4LayerHostWeights& L1 = w.host.layers[1]; + CHECK(L1.comp_ape.size() == static_cast(cr * 2 * hd)); + CHECK(L1.comp_wgate.size() == static_cast(2 * hd * H)); + CHECK(L1.idx_wk.size() == static_cast(2 * ihd * H)); + CHECK(L1.idx_wq.size() == static_cast(inh * ihd * w.params.q_lora_rank)); + // The two upstream does NOT widen (`compressor.py:288` is + // `RMSNorm(self.head_dim, self.rms_norm_eps)`, the file's only `RMSNorm(`; + // `weights_proj` is `[n_head, hidden_size]`). + CHECK(L1.comp_norm_weight.size() == static_cast(hd)); + CHECK(L1.idx_wproj.size() == static_cast(inh * H)); + + // 2. AND THE FORWARD REFUSES BY NAME. Be exact about what that is worth: + // `Gemm`'s host arm is a `MatVec` whose size assertion is UNCONDITIONAL + // (`deepseek_v4.cpp:413`, a `VT_CHECK` throw and not an `assert`), so + // without this the wide `comp_wgate` does NOT emit a plausible wrong + // number — it throws `vt: MatVec weight size mismatch at + // deepseek_v4.cpp:413`, naming no tensor, no layer and nothing missing. + // What this case gates is the DIAGNOSTIC: that the refusal arrives instead, + // and carries the layer, every mismatched tensor, both counts and the + // missing capability. The `!msg.empty()` assertion below therefore does not + // on its own separate the two, which is why every `Mentions` check that + // follows it is part of the gate and not decoration. + QueueGuard g; + const vllm::v1::CommonAttentionMetadata meta{}; + const std::vector kv; + const std::string msg = dsv4_exl3_fixture::ThrowMessage([&] { + (void)vllm::DeepseekV4Model::Forward(kTokens, kPositions, meta, kv, w, g.q, {}); + }); + CAPTURE(msg); + REQUIRE(!msg.empty()); // a forward that RETURNS here mis-indexed and did not say so + + // The AFFECTED LAYER, by number. + CHECK(dsv4_exl3_fixture::Mentions(msg, "layer 1")); + + // EVERY tensor whose width it cannot index, each with both counts. All four + // are named in ONE message on purpose: a refusal that stopped at the first + // mismatch would make the other three checks unfalsifiable, because deleting + // any one of them would still red on the first. + CHECK(dsv4_exl3_fixture::Mentions(msg, "compressor.ape")); + CHECK(dsv4_exl3_fixture::Mentions(msg, "compressor.wgate.weight")); + CHECK(dsv4_exl3_fixture::Mentions(msg, "indexer.compressor.wkv.weight")); + CHECK(dsv4_exl3_fixture::Mentions(msg, "indexer.wq_b")); + // Both counts, BOUND to the tensor that reported them. Asserting the numbers + // on their own does not gate this: several coincide at these dimensions (see + // `MismatchLine`), so a wrong count on one tensor is still found on another's + // line. `indexer.wq_b` gets its counts here too — it had none before, and its + // indexed count is one of the colliding 2048s. + CHECK(dsv4_exl3_fixture::Mentions( + msg, MismatchLine("compressor.ape", "[compress_ratio, head_dim]", cr * hd, + cr * 2 * hd))); + CHECK(dsv4_exl3_fixture::Mentions( + msg, MismatchLine("compressor.wgate.weight", "[head_dim, hidden_size]", hd * H, + 2 * hd * H))); + CHECK(dsv4_exl3_fixture::Mentions( + msg, MismatchLine("indexer.compressor.wkv.weight", "[index_head_dim, hidden_size]", + ihd * H, 2 * ihd * H))); + CHECK(dsv4_exl3_fixture::Mentions( + msg, MismatchLine("indexer.wq_b", "[index_n_heads*index_head_dim, hidden_size]", + inh * ihd * H, inh * ihd * w.params.q_lora_rank))); + + // The MISSING CAPABILITY, named — the AGENTS.md rule this case exists for. + CHECK(dsv4_exl3_fixture::Mentions(msg, "coff")); + CHECK(dsv4_exl3_fixture::Mentions(msg, "compressor.py:247-248")); + CHECK(dsv4_exl3_fixture::Mentions(msg, "q_lora_rank")); + CHECK(dsv4_exl3_fixture::Mentions(msg, "1970")); + + // 3. AND A COMPRESSOR LAYER THE FORWARD CAN INDEX STILL RUNS. The refusal is + // keyed on the width the forward indexes, not on the weight SOURCE — which + // is the mistake `dsa_dense = (be.gguf != nullptr)` makes (#1964). Without + // this case a check that simply refused every EXL3 compressor layer would + // pass everything above. `ForwardFixtureOptions()` is `cr == 128`, where + // `coff` is 1 and the loaded width IS the indexed one. + auto fc = BuildFixture(ForwardFixtureOptions()); + const DeepseekV4Weights wc = + vllm::LoadDeepseekV4ForCausalLMWeights(fc->shards, fc->config); + const std::vector logits = + vllm::DeepseekV4Model::Forward(kTokens, kPositions, meta, kv, wc, g.q, {}); + CHECK(AllFinite(logits)); + + // 4. THE TWO CHECKS THE LOADER MAKES UNREACHABLE ARE STILL FALSIFIABLE. + // `comp_norm_weight` and `idx_wproj` are the two DSA slots upstream does + // NOT widen (`compressor.py:288`; `weights_proj` is `[n_head, hidden_size]`), + // so the loader requires them at exactly the width the forward indexes and + // NO checkpoint can reach their checks. Deleting either one leaves both + // suites green, which would leave "all six are load-bearing" as an + // impression the gate does not support. + // + // So they are gated by MUTATING a tower the production loader produced, + // rather than by a checkpoint. Be exact about what that proves and what it + // does not: it proves the two checks FIRE and are named in the message, and + // it proves nothing about reachability, because the state it constructs is + // one the loader cannot emit. They are defensive checks against a future + // loader/forward disagreement, and this is the falsifiability half only. + // + // They are mutated ONE AT A TIME, in separate forwards. Mutating both at + // once and asserting on one message cannot separate them here: `head_dim - + // 1` and `index_n_heads * hidden_size - 1` are the SAME number at this + // fixture (511), so the two count assertions were one assertion written + // twice. One tensor per message makes each count unambiguous. + auto RefusalAfter = [&](void (*mutate)(DeepseekV4LayerHostWeights&)) { + DeepseekV4Weights wm = w; // the REAL-geometry tower, which already refuses + mutate(wm.host.layers[1]); + const std::string m = dsv4_exl3_fixture::ThrowMessage([&] { + (void)vllm::DeepseekV4Model::Forward(kTokens, kPositions, meta, kv, wm, g.q, {}); + }); + REQUIRE(!m.empty()); + return m; + }; + + const std::string nmsg = + RefusalAfter([](DeepseekV4LayerHostWeights& L) { L.comp_norm_weight.pop_back(); }); + CAPTURE(nmsg); + CHECK(dsv4_exl3_fixture::Mentions(nmsg, "compressor.norm.weight")); + CHECK(dsv4_exl3_fixture::Mentions( + nmsg, MismatchLine("compressor.norm.weight", "[head_dim]", hd, hd - 1))); + + const std::string pmsg = + RefusalAfter([](DeepseekV4LayerHostWeights& L) { L.idx_wproj.pop_back(); }); + CAPTURE(pmsg); + CHECK(dsv4_exl3_fixture::Mentions(pmsg, "weights_proj.weight")); + CHECK(dsv4_exl3_fixture::Mentions( + pmsg, MismatchLine("indexer.weights_proj.weight", "[index_n_heads, hidden_size]", + inh * H, inh * H - 1))); + // ...and each mutation names ONLY its own tensor, which is what makes the two + // checks separately falsifiable rather than jointly. + CHECK(!dsv4_exl3_fixture::Mentions(nmsg, "weights_proj.weight")); + CHECK(!dsv4_exl3_fixture::Mentions(pmsg, "compressor.norm.weight")); +} diff --git a/tests/vllm/models/test_deepseek_v4_exl3_loader.cpp b/tests/vllm/models/test_deepseek_v4_exl3_loader.cpp index 60de854fa..5de57a231 100644 --- a/tests/vllm/models/test_deepseek_v4_exl3_loader.cpp +++ b/tests/vllm/models/test_deepseek_v4_exl3_loader.cpp @@ -376,6 +376,15 @@ TEST_CASE("dsv4 exl3 W1c: the load materializes the carried tower and sets the f opt.index_n_heads = 2; opt.index_head_dim = 4; opt.index_topk = 3; + // The REAL artifact's geometry, which is the only one upstream emits at + // `cr == 4`: `coff = 1 + (compress_ratio == 4)` is 2 + // (`vllm/models/deepseek_v4/compressor.py:247-248`), so the compressor and + // indexer-compressor families are doubled while both norms and `weights_proj` + // are not, and `indexer.wq_b`'s K is `q_lora_rank` (`attention.py:721-726`). + // This is the tower the loader must materialize; whether the host forward can + // INDEX it is a separate question, answered by the refusal gated in + // `test_deepseek_v4_exl3_forward.cpp` (#1970). + opt.real_dsa_geometry = true; auto f = BuildFixture(opt); const vllm::DeepseekV4Weights w = vllm::LoadDeepseekV4ForCausalLMWeights(f->shards, f->config); @@ -411,17 +420,22 @@ TEST_CASE("dsv4 exl3 W1c: the load materializes the carried tower and sets the f CHECK(L0.tid2eid.size() == static_cast(kVocab * opt.topk)); CHECK(L0.gate_bias.empty()); - // Layer 1 carries the DSA compressor + Lightning-Indexer at the COLLAPSED - // geometry the host forward indexes. + // Layer 1 carries the DSA compressor + Lightning-Indexer at the width UPSTREAM + // DERIVES for `cr == 4`. The doubled and undoubled members are asserted side + // by side on purpose: a loader that widened the whole family would pass an + // assertion that only checked the four doubled ones. const vllm::DeepseekV4LayerHostWeights& L1 = w.host.layers[1]; CHECK(L1.tid2eid.empty()); CHECK(L1.gate_bias.size() == static_cast(kExperts)); - CHECK(L1.comp_wgate.size() == static_cast(kHeadDim * H)); - CHECK(L1.comp_ape.size() == static_cast(4 * kHeadDim)); - CHECK(L1.comp_norm_weight.size() == static_cast(kHeadDim)); + CHECK(L1.comp_wgate.size() == static_cast(2 * kHeadDim * H)); + CHECK(L1.comp_ape.size() == static_cast(4 * 2 * kHeadDim)); + CHECK(L1.idx_wk.size() == static_cast(2 * opt.index_head_dim * H)); CHECK(L1.idx_wq.size() == - static_cast(opt.index_n_heads * opt.index_head_dim * H)); - CHECK(L1.idx_wk.size() == static_cast(opt.index_head_dim * H)); + static_cast(opt.index_n_heads * opt.index_head_dim * kQLora)); + // NOT widened by `coff`: the compressor norm (`compressor.py:288` is + // `RMSNorm(self.head_dim, self.rms_norm_eps)`) and `weights_proj`, which is + // `[n_head, hidden_size]`. + CHECK(L1.comp_norm_weight.size() == static_cast(kHeadDim)); CHECK(L1.idx_wproj.size() == static_cast(opt.index_n_heads * H)); // The routed experts are the TRELLIS tower; a second host copy of them would @@ -598,33 +612,152 @@ TEST_CASE("dsv4 exl3 W1c: the carried tower is read from MISALIGNED payloads") { } TEST_CASE("dsv4 exl3 W1c: a carried tensor this arm cannot route REFUSES BY NAME") { - SUBCASE("the REAL artifact's 2*head_dim compressor") { - // MEASURED on `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32` - // (2026-08-25): `layers.N.attn.compressor.wgate.weight` is BF16 - // [1024, 4096] = [2*head_dim, H], the ds4 `coff = 2` width, while the host - // forward's compressor indexes [head_dim, H]. `Gemm`'s host arm is a - // `MatVec` with no length check, so materializing the wide tensor into that - // slot is a SILENTLY WRONG number rather than a crash — the refusal is what - // keeps it from being one. 41 of the real artifact's 43 layers carry a - // compressor, so this is the shape that stops the real checkpoint, and the - // spec's `## Owed` names what would close it. + // RETIRED by #1970: this subcase asserted that the loader REFUSES the real + // artifact's `coff = 2` compressor, which is exactly the behaviour option C + // removes. The refusal did not disappear — it moved to the forward, where the + // mis-index it prevents actually lives, and it is re-gated end to end by + // `test_deepseek_v4_exl3_forward.cpp` ("the REAL DSA geometry LOADS and the + // FORWARD refuses by name"), which drives the production loader AND the + // production forward rather than the loader alone. + // + // Its fixture was also wrong about the artifact: it set `compress_ratios = + // {128}` while doubling the width, and upstream's `coff` is + // `1 + (compress_ratio == 4)` (`vllm/models/deepseek_v4/compressor.py:247-248`), + // so a `cr == 128` layer is NOT doubled and that shape is one the checkpoint + // does not store. The fixture flag it used is now `real_dsa_geometry` and + // applies the per-layer rule instead. + SUBCASE("a DSA width that is NEITHER geometry") { + // #1970 moved the compressor family's width from a CONSTANT to upstream's + // own derivation, `coff = 1 + (compress_ratio == 4)` + // (`compressor.py:247-248`). It did not stop checking, and a fixture flag + // that only toggled between the derived width and the collapsed one could + // not tell a derivation from a two-value allow-list. This writes a third + // width no oracle emits. FixtureOptions opt; opt.layers = 1; opt.compress_ratios = {128}; - opt.real_compressor_width = true; + opt.bogus_dsa_width = true; + auto f = BuildFixture(opt); + const std::string msg = ThrowMessage( + [&] { vllm::LoadDeepseekV4ForCausalLMWeights(f->shards, f->config); }); + CAPTURE(msg); + CHECK(Mentions(msg, "compressor.wgate.weight")); + CHECK(Mentions(msg, "MODEL-DSV4-EXL3")); + CHECK(Mentions(msg, "W1c")); + // The DERIVED width is named, so the message says what WOULD be routable + // rather than only that this one is not. At cr == 128 `coff` is 1. + CHECK(Mentions(msg, "512")); + CHECK(Mentions(msg, "1536")); // the refused 3 * head_dim + } + SUBCASE("a COLLAPSED `cr == 4` family, which upstream cannot emit") { + // The width is DERIVED from `coff = 1 + (compress_ratio == 4)` + // (`vllm/models/deepseek_v4/compressor.py:247-248` at the parity pin + // `5559679229bc961848b121ccdeaa8fa5d79bec98`), so at `cr == 4` upstream + // emits ONE width and it is the doubled one. A `cr == 4` checkpoint whose + // compressor family is UNDOUBLED is a checkpoint upstream cannot load at + // all, and this loader must not accept it either. + // + // This is the case a two-width loader accepted. It is separated from the + // "neither geometry" subcase on purpose: a third width proves the loader + // still checks SOMETHING, and only this one proves it checks the width + // upstream actually derives rather than a set chosen to suit a fixture. + FixtureOptions opt; + opt.layers = 1; + opt.compress_ratios = {4}; + opt.index_n_heads = 2; + opt.index_head_dim = 4; + opt.index_topk = 3; + opt.real_dsa_geometry = false; // the collapsed, undoubled family + auto f = BuildFixture(opt); + const std::string msg = ThrowMessage( + [&] { vllm::LoadDeepseekV4ForCausalLMWeights(f->shards, f->config); }); + CAPTURE(msg); + CHECK(Mentions(msg, "compressor.wgate.weight")); + CHECK(Mentions(msg, "coff")); + CHECK(Mentions(msg, "compressor.py:247-248")); + CHECK(Mentions(msg, "MODEL-DSV4-EXL3")); + CHECK(Mentions(msg, "W1c")); + CHECK(Mentions(msg, "1024")); // the derived `coff * head_dim` + CHECK(Mentions(msg, "512")); // the collapsed width the checkpoint carries + } + SUBCASE("`indexer.wq_b` at the WRONG INPUT SPACE") { + // `wq_b` is the one DSA tensor whose width is NOT a `coff` width. Upstream + // builds it as `ReplicatedLinear(q_lora_rank, head_dim * n_head)` + // (`attention.py:721-726`) and calls it on `qr`, the q-LoRA latent + // (`:835`), so its K is `q_lora_rank` and nothing about it is doubled. + // + // It gets its own case because the rest of the family refuses FIRST: a + // wholly collapsed checkpoint reds on `compressor.wgate.weight` and never + // reaches this check, which would leave the loader free to accept a `wq_b` + // at `hidden_size` unnoticed. So this writes the real geometry everywhere + // ELSE and collapses only this one tensor. + FixtureOptions opt; + opt.layers = 1; + opt.compress_ratios = {4}; + opt.index_n_heads = 2; + opt.index_head_dim = 4; + opt.index_topk = 3; + opt.real_dsa_geometry = true; + opt.collapsed_indexer_wq_b = true; + auto f = BuildFixture(opt); + const std::string msg = ThrowMessage( + [&] { vllm::LoadDeepseekV4ForCausalLMWeights(f->shards, f->config); }); + CAPTURE(msg); + CHECK(Mentions(msg, "indexer.wq_b")); + CHECK(Mentions(msg, "q_lora_rank")); + CHECK(Mentions(msg, "attention.py:721-726")); + // The message must NOT blame `coff`, which has nothing to do with this K. + CHECK(!Mentions(msg, "compressor.py:247-248")); + CHECK(Mentions(msg, std::to_string(kQLora))); // 128, the required K + CHECK(Mentions(msg, std::to_string(kHidden))); // 256, the K carried + } + SUBCASE("the INDEXER's own compressor at the COLLAPSED width") { + // The third derivation. `coff` widens the indexer's own + // `DeepseekCompressor` family too, because upstream builds it at + // `head_dim = index_head_dim` with the SAME ratio + // (`vllm/models/deepseek_v4/attention.py:768-776`, the `DeepseekV4Indexer` + // one — `attention.py:335` is the attention's own) and `coff` is a property + // of the ratio, not of the head width. So `indexer.compressor.wkv.weight` + // is derived at `coff * index_head_dim` by a rule that is the SAME as the + // main compressor's but reads a different head width. + // + // It gets its own case for the reason `collapsed_indexer_wq_b` does, and + // the round-2 fresh review found the hole by mutation: with the whole + // family collapsed the MAIN compressor refuses first, so this derivation's + // message was never read and DELETING the check left both suites green. + // Here everything else is at the real geometry and only this tensor is + // collapsed, so the loader must refuse BY NAME on this one. + FixtureOptions opt; + opt.layers = 1; + opt.compress_ratios = {4}; + opt.index_n_heads = 2; + opt.index_head_dim = 4; + opt.index_topk = 3; + opt.real_dsa_geometry = true; + opt.collapsed_indexer_wkv = true; auto f = BuildFixture(opt); const std::string msg = ThrowMessage( [&] { vllm::LoadDeepseekV4ForCausalLMWeights(f->shards, f->config); }); CAPTURE(msg); - CHECK(Mentions(msg, "compressor")); + CHECK(Mentions(msg, "indexer.compressor.wkv.weight")); + // Its OWN derivation, named: `coff` over `index_head_dim`, and WHY `coff` + // is 2 here at all. A message that named only `coff * head_dim` would be + // describing the main compressor's rule. + CHECK(Mentions(msg, "coff * index_head_dim")); + CHECK(Mentions(msg, "compress_ratio == 4")); + CHECK(Mentions(msg, "attention.py:274")); + CHECK(Mentions(msg, "compressor.py:247-248")); CHECK(Mentions(msg, "MODEL-DSV4-EXL3")); CHECK(Mentions(msg, "W1c")); - // The REASON, not just the fact of a throw: both shapes in the message. - // `compressor.ape` is the first of the family the loader reaches, so it is - // the one that names the width — [compress_ratio, head_dim] wanted against - // the artifact's [compress_ratio, 2*head_dim]. - CHECK(Mentions(msg, "[128,512]")); - CHECK(Mentions(msg, "[128,1024]")); + // Both widths, bound to this tensor: 8 = coff * index_head_dim is what + // upstream derives, 4 = index_head_dim is what this checkpoint carries. + CHECK(Mentions(msg, "dimension 0 must be 8")); + CHECK(Mentions(msg, "got 4 in [4,256]")); + // The main compressor is at the REAL width here, so it must NOT be what + // refused — otherwise this case would pass without ever reaching the + // derivation it exists to gate. + CHECK(!Mentions(msg, "compressor.wgate.weight")); + CHECK(!Mentions(msg, "indexer.wq_b")); } SUBCASE("no recipe for the carried FP8 half") { // The carried MLA linears are block-wise FP8 and the block size comes from