diff --git a/.agents/issue-index.md b/.agents/issue-index.md index adaa0c8c9..6ed3766c5 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -748,6 +748,8 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#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 | | [#1978](https://github.com/mudler/vllm.cpp/issues/1978) | `MODEL-MM-QWEN4-EXP` | **`Qwen/Qwen3.8-Flash-Next` declares `Qwen4ExpForConditionalGeneration` / `qwen4_exp`, a new architecture vLLM does not implement, so the port runs on a split oracle: transformers for the ALGORITHM, vLLM ops for the OPTIMIZED PATH.** Released 2026-08-24, 180B total / 6B activated, image-text-to-text. The `Qwen3.8` in the name is marketing continuity: `.agents/specs/qwen38-27b-bf16-gate.md`'s "one config key differs" precedent does NOT extend here. Read live 2026-08-26 at vLLM `origin/main` = `6a5e8f5979`: no `qwen4*` path, no registry entry, and a repository-wide GitHub search for `qwen4` returns ZERO results; `vllm-omni` likewise. That is absence from vLLM `main` rather than staleness in our pin `555967922`, so a pin advance does not reach it. What exists is transformers [#48337](https://github.com/huggingface/transformers/pull/48337) "Add Qwen4Exp model", MERGED 2026-08-26, 5211 lines, and SGLang [#36497](https://github.com/sgl-project/sglang/pull/36497), still OPEN and therefore inadmissible. **Developer direction 2026-08-26, recorded verbatim: "use transformers as oracle for algorithmic side. but use ops from vllm so we account for optimized path."** Justified rather than convenient: `Qwen4ExpTextQSAIndexer.forward` loops in Python over `(batch_idx, query_idx)` and says "we only allow eager and sdpa", so porting it as written yields a correct model at an indefensible speed, while AGENTS.md's mirror-vLLM polarity still binds every primitive vLLM implements. `Qwen4ExpTextModel` inherits from `Qwen3_5MoeTextModel` and leaves rotary, MLP, experts, TopK router and the ENTIRE vision tower unchanged (`class Qwen4ExpVisionModel(Qwen3_5MoeVisionModel): pass`), all of which this tree has; GDN is an exact match for our AOT gate (`K=V=128, Hg=16, Hv=48` against `src/vt/cuda/cuda_gdn.cu`'s `H in {48,32}`). The delta is four things, and **exactly two have no vLLM op at all**: the PLE dilated depthwise conv (kernel 4, dilation 3; `git grep dilation` over vLLM `layers/mamba/` = 0 hits) and the n-gram hashed embedding. **The survey's load-bearing finding, and it REVERSES this row's first reading: QSA's structural twin is DeepSeek-V4's C4 indexer lane, NOT MiniMax-M3.** The original call was that QSA, being plain GQA rather than MLA, had to map onto vLLM's non-MLA block-sparse case; that reasoning rested on treating `MLAAttentionSpec` as an MLA claim, and **it is not one** — M3's own indexer cache uses it while M3 is a plain-GQA model, with the comment "Key-only: MLAAttentionSpec budgets one vector/token (not 2x for K+V)". It is a budget shape. Remove that prop and the GQA-vs-MLA argument collapses. Verified line by line at `6a5e8f5979`: **nine independent structural matches with DSv4**, `compress_ratio == 4` literally the same number — MQA index with 1 key head at dim 128; `relu(q.k)` summed over index heads vs `(score.relu() * weights).sum(dim=0)`; `1/sqrt(head_dim)`; one score set per query token with no head axis vs `topk_indices_buffer[num_tokens, topk]`; pooling boundary `(position+1) % COMPRESS_RATIO == 0`; RMSNorm on the pooled key; **RoPE at the block-start position** vs `compressed_pos = (position // CR) * CR`; candidate count `visible // compress_ratio`; and one stored state per 4 tokens via `MLAAttentionSpec(tokens_per_state=compress_ratio)`, a first-class KV field documented as "Ints > 1 compress multiple tokens into one state (DSv4 sparse MLA)" that has no M3 equivalent. **M3 is a DIFFERENT ALGORITHM**, not a worse fit: its score is `tl.max(qk, axis=1)` over 128 RAW token dots with no pooling, no relu and no head reduction, it asserts `num_idx_heads == num_kv_heads` ("no topk index reduce") so it emits one block set PER KV HEAD, and its `SPARSE_BLOCK_SIZE = 128` is welded to the KV page size ("One sparse block == one KV page") on both the score and the attend side — moving it to 4 forces a page size of 4 and breaks `tl.dot`, whose tile needs >= 16. M3 contributes exactly ONE thing and it is a wiring precedent, not an algorithm: that a plain-GQA model can own a key-only side cache through `MLAAttentionSpec` and a private indexer backend. **The genuinely new work is the CONSUMER and nothing upstream supplies it** — every DSv4 sparse consumer attends to COMPRESSED MLA KV (one state per 4 tokens) and M3's attend to raw tokens only at page granularity, while QSA attends to RAW tokens selected at ratio-4 granularity. Two silent-failure traps follow: wiring QSA's top-k into a DSv4 sparse-MLA consumer attends a POOLED key/value and still emits plausible tokens, and **a short-prompt token gate cannot catch it because at context <= `indexer_budget` 2048 every candidate is selected** — so any QSA gate must run past 2048 tokens of context, which is now a stated `## Gates` requirement; and `SparseAttnCompressNormRopeStoreC4Kernel` does NOT mean-pool despite its name — it is a learned softmax pool over an OVERLAPPING window of 8 using a score channel this checkpoint does not have, and the CuteDSL variant refuses `overlap=False` at compile, so the **Triton** `head_dim=128` variant is the correct starting point. Two structural consequences beyond the module list: the residual stream is `hc_count * hidden_size` = **4 x 2560 = 10240 wide through the whole stack** with a low-rank read gate and per-branch scalar write gate around both attention and MLP, which is a change to the per-layer loop and every residual buffer rather than a drop-in module; and `number_of_conv_states = 3` on a PLE layer (GDN conv, PLE conv, n-gram token history) plus the indexer side cache, adjacent to [#1963](https://github.com/mudler/vllm.cpp/issues/1963) and [#1966](https://github.com/mudler/vllm.cpp/issues/1966). **NOTHING PUBLISHED FITS**, read live from the HF API against ~119 GB usable on GB10: BF16 ~360 GB, official FP8 ~180 GB, `RadixArk/...-NVFP4` ~128 GB (NVFP4 backbone with the n-gram table left at FP8, 51.2 GB) and `unsloth/...-GGUF` is a README with ZERO weight files. No GGUF exists and no tool can make one, because llama.cpp has no `qwen4_exp` either, so the standing k-quant requirement means authoring the arch on our side AND states that the quantized arms have NO llama.cpp oracle. **The chosen arm does NOT load today, and the blocker is neither the offload nor the budget: this tree cannot keep a gather table quantized at all.** `KeepQuantKDim` returns `-1` for `GgufTensorRole::kEmbeddingTable` (`src/vllm/model_executor/model_loader/gguf_keep_quant.cpp`), and `qwen3_5_gguf_weights.cpp` asserts it by name — "the embedding table cannot keep quant blocks" — so a Q4_K or Q8_0 n-gram table EXPANDS to bf16 and 51.2B params become **102.4 GB of anonymous memory**; the arm dies before the first forward. The reason was already sitting in a header comment ("a gather, not a GEMM ... A quantized-gather op is a follow-up row") and **no such row exists**. The only non-expanding gather residency is `kKeepF16`, requiring ggml type 1 exactly (102.4 GB on disk) and CPU-ONLY, because `EmbeddingKernelCuda` refuses anything but f32/bf16. **Second blocker:** `moe_intermediate_size = 640` makes `ffn_down_exps` Q4_K-illegal on its reduction dim (640 % 256 = 128), as does `hc_lowrank = 320`; llama.cpp's substitution is believed to be Q5_0 (**UNVERIFIED, owed against the pinned llama.cpp oracle**) and the dependent fact IS verified in-tree — our reader knows ggml ids `0,1,2,8,10..14,16,18,19,22..28,30,39,40,41,66` and has **no entry for 3, 6, 7 or 20**, so a stock `llama-quantize -Q4_K_M` file fails at header parse. We author the converter, so the fix is Q4_0 (block 32, same 4.5 bpw). **`ENG-WEIGHT-OFFLOAD` will not help** — it moves zero bytes today (`ConsiderWeight` has no production callers, pinned by a test) and is documented inert on GB10; the tier that DOES work already ships and is proven by the 2.4T model serving 369.97 GiB from a 119.631 GiB box at ~62 GiB resident: mmap `MAP_PRIVATE`, borrow in place, alias the host pointer, `prefault: false`. Corrected sizing: backbone ~67.7 GiB, whole process ~73.5 GiB of 119.631 at 32K single-stream, ~46 GiB of headroom for the page cache, so the ~76 GB estimate was right within 10%. The design works because per-token demand is **<= 64 KiB of reads** (16 lookups x 160 dims over at most 16 pages) against the 2.4T expert lane's 6.95 GB/token. The architecture supplies its own lever: the per-token n-gram cost is `(ngram_size-1)*heads_per_ngram` = 16 lookups of 160 dims, so **51 GB of the 180 GB, 28% of the model, is a table touched 16 times per token** and making it non-resident is the intended design point (RadixArk reached the same split independently). Sizing arithmetic, NOT measurement: Q8_0 throughout ~191 GB (no), Q4_K_M throughout ~109 GB (yes, ~10 GB left for KV), Q4_K_M backbone with the table non-resident ~76 GB. GB10 is UNIFIED memory so "offload to host" is not a move there; non-resident means disk-backed, and its cost is unmeasured. **Two decisions were put to the developer as explicit accept-or-reject and BOTH are settled 2026-08-26, recorded in place rather than left open.** (1) `.agents/oracles/transformers.md` pins transformers to 5.14.1, deliberately tied to what the pinned vLLM environment resolves so the environment cannot hold two `transformers` at once, and **5.14.1 does not contain `Qwen4Exp`**; the lane-scoped second pin is **ACCEPTED**, on the argument that the invariant guards a vLLM environment against drifting from its transformers and here there is no vLLM implementation to drift from, and it expires the moment vLLM registers `qwen4_exp`. **The lane pin is a real release, not a branch SHA**, which was not the expected outcome: `Qwen4Exp` merged to `main` at 12:03:40Z on 2026-08-26 and `v5.16.0` published at 12:35:15Z, and this was BOUNDED rather than assumed by fetching `models/qwen4_exp/modeling_qwen4_exp.py` at each tag — `v5.16.0` HTTP **200**, `v5.15.0` HTTP **404** — making 5.16.0 the FIRST release carrying the architecture and therefore the tightest available pin. The version string is UNMEASURED (it is the release proven to contain the model, not a `transformers.__version__` read off a running oracle) and `gateable` stays `no`. (2) The first runnable arm is the **Q4_K_M backbone with the n-gram table NON-RESIDENT** (~76 GB). Q8_0 was raised and does not fit at ~191 GB, and no partial-Q8 split reaches 119 GB with the backbone at 8 bits; Q4_K_M-throughout fits on paper at ~109 GB but leaves ~10 GB for KV and activations on a 262144-native-context model, which is not a margin. This promotes the non-resident table from a note to a first-class W6 deliverable, and it is NOT free: GB10 is UNIFIED memory, so the existing host-pinned offload seam (`ENG-WEIGHT-OFFLOAD`, mirroring vLLM's `cpu_offload_gb`) does not by itself solve it there and the mechanism must be disk-backed or genuinely unloaded — established before it is designed around. Spec: [`specs/qwen4-exp-flash-next.md`](specs/qwen4-exp-flash-next.md). No product code lands under the spec pull request | feature | +| [#2008](https://github.com/mudler/vllm.cpp/issues/2008) | `SPEC-DFLASH2` | **DFlash2 serves exactly one sequence: the draft context is keyed by BATCH ROW, and `InputBatch::condense` moves a live request between rows.** Measured on an idle leased GB10 at c=1 (24.70 out tok/s, TPOT 37.90 ms, 8/8 ok) and c=2 (VOID, ok=1 failed=7, `propose_drafts_block: context position discontinuity`, then every later request `[request submitted to a stopped AsyncLLM]`). The operator's isolation settles the layer: with `--speculative-config` omitted and everything else identical, both concurrent requests complete, so batching, scheduling, the paged KV cache, the block tables and the sampler are all correct. `GPUModelRunner`'s four draft arrays (`include/vllm/v1/worker/gpu/runner.h:852-870`) are indexed by row; `condense` slides a live request into a departed neighbour's row (`src/vllm/v1/worker/gpu/input_batch.cpp:686-706`) and `swap_states` exchanges two live rows (`:762-847`), permuting every per-slot array they own — including the block-table rows — but knowing nothing about the runner's four. The survivor then meets the departed request's bookkeeping, the reuse test at `runner.cpp:2895-2906` resets its store to empty, and the invariant at `:2939-2945` correctly refuses rather than drafting from a foreign context. `ok=1` is the mechanism's signature, not an incidental count. Upstream has no analogue of the host-side counter at all: read beyond-pin at `b389ac2946`, DFlash/DFlash2 address the draft KV by ABSOLUTE POSITION (`dflash/speculator.py:562-590`), re-read the anchor from the target each step (`:553`), and index every cross-step tensor by the persistent request slot via `idx_mapping` (`:536`, `dflash2/speculator.py:95`) in a V2 runner that has no `condense` at all (`gpu/states.py:29,100,132`); the legacy V1 runner does condense and carries the draft's block-table row with the request (`gpu_input_batch.py:786`). Fixed by keying the four arrays on request id. Two things are OWED and named in the spec rather than folded in: the `P == 1` capture gate at `src/vllm/model_executor/models/qwen3_dflash.cpp:1577`, which this is the first change to make measurable because no batch previously survived to `P > 1`; and a distinct c=1 defect this found and did not fix — a prefix-cache hit or a resumed request is admitted with `num_computed_tokens > 0`, has no draft context for the cache-supplied tokens, and trips the same invariant, which the #2008 measurement never met because it ran `--no-enable-prefix-caching`. Spec: [`specs/dflash2-request-scoped-context.md`](specs/dflash2-request-scoped-context.md) | bug | +| [#2009](https://github.com/mudler/vllm.cpp/issues/2009) | — | **DFlash2's draft-context position invariant is ungated: deleting it leaves the suite green.** Found while fixing [#2008](https://github.com/mudler/vllm.cpp/issues/2008) and owed by [`specs/dflash2-request-scoped-context.md`](specs/dflash2-request-scoped-context.md). `src/vllm/v1/worker/gpu/runner.cpp:2939-2945` is the guard the whole draft-context accumulation rests on and the reason #2008 was a loud refusal rather than a silent wrong-context draft. Measured: deleted on the pre-#2008 code, `test_dflash2_concurrency` stays green at 2 cases / 30 assertions, the row move resets the survivor's store to empty, it drafts from a context that is not its own, and nothing notices — because the verify is lossless, so a draft from the wrong context costs acceptance and never a token, and every token-shaped gate in this tree is blind to it by construction. #2008's own gate cannot close this: the natural leg, comparing drafted blocks against a solo control, is a **tautology** on the shared DFlash2 fixture — with the invariant deleted and the context reset at every row move the draft still emits `12 12 12` at every step of both runs, because its seeded-noise weights over a 24-token vocabulary collapse the selector walk to one id, so nine passing string comparisons measured nothing. That leg was written, run and removed rather than shipped. `test_dflash2_runner_reach`'s value-sensitivity case is unaffected — it moves the drafts by changing the selector's WEIGHTS, not the context. Closing this needs a fixture whose drafted block is demonstrably sensitive to the draft CONTEXT, which is a fixture problem before it is a test problem and is the same instrument several DFlash2 rows would benefit from: today the tree can prove a draft moves with its weights and cannot prove it moves with its context | bug | | [#1981](https://github.com/mudler/vllm.cpp/issues/1981) | `MODEL-MM-QWEN4-EXP` | **W1 of [#1978](https://github.com/mudler/vllm.cpp/issues/1978): the `qwen4_exp` config surface — resolve, validate, register, and refuse by name everywhere else.** Filed and closed in flow. It is indexed rather than left to the pull request body because every `Refuse()` message this code emits ends "See `.agents/specs/qwen4-exp-flash-next.md` and issue #1981", so a reader who follows the pointer a running binary gives them has to find the issue at the other end of it; AGENTS.md requires the index, the spec and the PR body to agree, and until this row only the PR body carried it. **The row's product is a BOUNDARY, and the boundary is measured.** `Qwen4ExpForConditionalGeneration` has no reachable token gate (`gateable = no`, nothing published fits a fleet device), so no downstream gate will ever catch a wrong config default by running the model, and the config layer is the last place one is checkable. The config layer itself IS gateable even though the model is not: `transformers` 5.16.0 installs and imports without torch and runs `validate_architecture` in full, so W1 is gated by a 39-case two-direction sweep — each config put through `Qwen4ExpConfig.from_dict` on one side and `LoadHfConfig -> ModelRegistry::Resolve -> factory->parse_config` on the other. **35 agree; 4 differ, and all 4 are ours refusing what upstream accepts**, never the reverse. All 15 upstream `validate_architecture` rejections are implemented and tabulated against their upstream line in the spec's `## The refusal boundary`, with the local tighter guards listed beside them. Four defaults were wrong in the first draft and every one of them is invisible to a token gate: `partial_rotary_factor` was read from the text config with a hardcoded 0.25 on the belief that `Qwen4ExpTextConfig` inherits it from `Qwen3_5MoeTextConfig` — the generated class is `class Qwen4ExpTextConfig(PreTrainedConfig)`, declares no such field, and `0.25` does not occur in the file, so the port both accepted configs upstream refuses (rotary_dim 64 where upstream computes 256 and raises) and refused one upstream accepts; the four PLE n-gram fields defaulted to 0 rather than 3 / 8 / 20000000 / 128, refusing a legal config and carrying a zero-sized n-gram vocabulary into W2; `output_gate_type` did not fall back to `hidden_act`, and its local check was a constant false the shared reader had already made unreachable; and `ple_embed_dim <= 0` was dropped from upstream's condition, so `-2560` passed the divisibility test because `-2560 % 16 == 0` in C++. Also landed: `eos_token_id` is now required when PLE is enabled (it is a segment boundary in the hashed n-gram construction, and the published GGUF stores it as `qwen4exp.ple.eos_token_id`); the forward refuses BEFORE the `ModelAs` downcast, because nothing can produce a loaded Qwen4-Exp while the loader refuses and a downcast placed first made the advertised refusal unreachable; `block_topk()` and `head_dim_per_ngram()` refuse instead of SIGFPE on a legally-parsed config with QSA or PLE absent; and the model's local `TextOf` now resolves `llm_config` and `thinker_config.text_config` like the shared `ResolveTextConfig`, which it did not, so one parse no longer answers "what is the text config" two different ways | bug | | [#1989](https://github.com/mudler/vllm.cpp/issues/1989) | `MODEL-MM-QWEN4-EXP` | **W6a: the GGUF reader had no `case 20`, so `GgufFile::OpenOne` died at header parse on shard 2 of the ONLY published Qwen3.8-Flash-Next artifact that fits any device this project owns.** `unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S` is 67.56 GiB in three shards against ~119.6 GiB usable on GB10, where bf16 is ~360 GB, the official FP8 ~180 GB and NVFP4 ~128 GB; the GGUF arm is therefore the path to a running model, not a follow-up to a safetensors one. Read live 2026-08-26 by HTTP range request over the shard headers: `general.architecture = "qwen4exp"`, `split.tensors.count = 1224` (shard 1 is 67 keys and ZERO tensors, shards 2 and 3 carry 595 and 629), `per_layer_token_embd.weight` IQ4_NL `[160, 320001536]`, `ffn_down_exps` IQ4_NL `[640, 2560, 512]`, `ffn_{gate,up}_exps` IQ1_S or IQ2_XXS, `indexer.{q,k}_proj` left BF16. IQ4_NL (id 20) appears 49 times and is unavoidable rather than a recipe preference: `moe_intermediate_size` 640 and the table row 160 are neither a multiple of 256, so no K-quant can encode them, and upstream's own `tensor_type_fallback` drops `IQ4_XS -> IQ4_NL` and `Q4_K -> Q5_0` — VERIFIED at the pin, `src/llama-quant.cpp:374-405 @ b10451`, discharging the spec's UNVERIFIED item; the same table maps `Q5_K -> Q5_1` (id 7), which we still lack, so a `-Q5_K_M` build of this model remains refused. **Second blocker, independent of the first:** a gather table could not be kept quantized AT ALL. `KeepQuantKDim` returned `-1` for `GgufTensorRole::kEmbeddingTable` and `qwen3_5_gguf_weights.cpp` asserted it by name, so a quantized 51.2 G-parameter n-gram table expanded to **102.4 GB of anonymous memory** against 28.8 GB of IQ4_NL blocks — the end of the box before the first forward. The `-1` was CORRECT until this row: without a dequantizing gather a kept table is bytes nothing can read. Both are closed. `vt::Embedding` now takes a block-quantized table and decodes ONE ROW per gathered id, a port of `ggml_compute_forward_get_rows_q` (`ggml/src/ggml-cpu/ops.cpp:4850 @ b10451`), and the table's residency follows the ordinary policy behind a gather-specific admission rule (`KeepQuantGatherDType`: a row DECODER, not the GEMM arm's `vec_dot`) and a device gate (`DeviceQuantGatherSupported`). Both new decoders are gated BIT-EXACTLY against the pinned llama.cpp decoding REAL bytes of the shipped tensor, read by range request at absolute offset 364622656 of shard 2 — ten IQ4_NL blocks, two whole gather rows — with the oracle built from a clean `git archive b10451` rather than from a working checkout. **The CUDA gather arm is OWED and it is the expensive half:** `EmbeddingKernelCuda` still refuses a block table, so on CUDA the table keeps expand-bf16, and a device-resident quantized table gathered on device is exactly the shape llama.cpp's #27742 does NOT have (it pins the n-gram table to the CPU by tensor class), which is where this model's high-concurrency advantage lives. Also landed: the `qwen4exp` config builder in its OWN translation unit with its own dispatch row, deliberately NOT reusing `HfConfigFromGguf`, which asserts its own three architectures by name and would refuse a fourth family as "qwen3_5 gguf:" — the #809 defect. Its key names follow llama.cpp #27742, which is what the shipped file uses, and it carries the architecture-specific numbers under the RELEASED `config.json`'s own spellings; `ple.layers` is the one exception, kept under its GGUF name because the file says `[1]` where config.json says `ple_layer_ids: [2]` and nothing in either resolves the offset. **Landed unreached, named per "Nothing lands dead":** the config builder IS reached through `kGgufArchArms`, but `ModelRegistry` does not resolve `Qwen4ExpForConditionalGeneration` — the model wiring is owed to [#1978](https://github.com/mudler/vllm.cpp/issues/1978) and listed under `## Owed` in `specs/qwen4-exp-flash-next.md`. No forward, no token claim and no speed claim from this wave | feature | | [#1988](https://github.com/mudler/vllm.cpp/issues/1988) | `MODEL-MM-QWEN4-EXP` | **W3 of the Qwen4-Exp port: the 4-branch GATED-RESIDUAL hyper-connection stream and the grouped RMSNorm it stands on.** The residual stream is `hc_count * hidden_size` = 4 x 2560 = 10240 wide through the whole 48-layer stack, read and written twice per layer, and collapsed at the end by the same class with its injection branch switched off — a change to every residual buffer, not a drop-in module. Landed here as a HOST reference (`src/vllm/model_executor/models/qwen4_exp_hc.{h,cpp}`) gated against goldens dumped by EXECUTING the lane-pinned oracle source: transformers `v5.16.0` `models/qwen4_exp/modeling_qwen4_exp.py` (sha256 `77fec77d…`), `Qwen4ExpTextRMSNorm` (:158-181) and `Qwen4ExpTextGatedResidual` (:941-969) lifted verbatim by line range, plus an independent double-precision reference. The grouped norm mirrors vLLM's op form — `RMSNormGated` (`layers/layernorm.py:172`, `group_size` at `:187`, grouped branch `:258-264`) with the gate disabled, NOT the plain `RMSNorm` (`:37`), whose only related knob is `var_hidden_size`, a prefix reduction that cannot express per-group norms. **Three findings the gate now pins.** (1) The `1 + w` parameterization: transformers applies `out * (1.0 + weight)` on a ZERO-init weight while vLLM applies `out * weight` on a ONES-init one, they coincide only under a load-time `w = 1.0 + w_hf`, and the published GGUF has that fold applied at CONVERT time — so it lives in exactly one named function, `HcNormWeightFromHf`, and skipping it scales every `hc_norm` by ~0 (reads as a checkpoint bug) while applying it twice scales by ~2x. (2) The two divisions by `hc_count` are different: one is INSIDE the SiLU on the `[320]` low-rank intermediate BEFORE the activation (`silu(down(x)/4)`, not `silu(down(x))/4`; SiLU is not homogeneous), the other is inside the injection sigmoid with the whole sigmoid scaled by 2 (`2*sigmoid(inject(x)/4)`, range (0,2), exactly 1.0 at a zero logit), and there is NO division on the up-projection sigmoid. (3) The elementwise multiply uses the NORMED stream, the reduce over hc is a MEAN and not a sum, and `hyper_input` is written back RAW. **The spec's `MhcPost`-with-identity-comb reuse claim is VERIFIED rather than trusted**, by a bit-equality case against our DeepSeek-V4 kernel; it holds on finite inputs and is not an identity for a negative-zero or non-finite residual, neither of which is reachable here. **Not reached at its merge commit** — W1 config registration (#1986) is still in review, so nothing loads a `qwen4_exp`; the wiring is owed by W5 (assembly) under #1978 and is listed in the spec's `## Owed`. No token claim and no speed claim: no arm of this model runs on any fleet device. | feature | diff --git a/.agents/specs/dflash2-request-scoped-context.md b/.agents/specs/dflash2-request-scoped-context.md new file mode 100644 index 000000000..5834d3430 --- /dev/null +++ b/.agents/specs/dflash2-request-scoped-context.md @@ -0,0 +1,427 @@ +# SPEC-DFLASH2 — the draft context belongs to the REQUEST, not to the batch ROW ([#2008](https://github.com/mudler/vllm.cpp/issues/2008)) + +Row: `SPEC-DFLASH2`. Issue: +[#2008](https://github.com/mudler/vllm.cpp/issues/2008). Parent waves: +[`dflash2-device-propose.md`](dflash2-device-propose.md) (W8, the device-resident +per-request store and the two invariants this defect trips) and +[`dflash2-ctx-store-capacity.md`](dflash2-ctx-store-capacity.md) (#1919, the +fallback branch a naive repair here would be mistaken for). + +## The finding + +Measured 2026-08-26 by the operator on `integ4/1574` @ `3d137890f`, artifact-gated +(`flash_fwd=1792`, `SpecDecodeFA2Bf16=1`, FA2 manifest `[121a]`), on an idle +leased GB10. DFlash2 K=8, `--num-blocks 3744 --max-num-seqs 16 --max-model-len +8192 --no-enable-prefix-caching`, `vllm bench serve --backend openai-chat`, +input 1024 / output 512: + +| rung | result | +|---|---| +| c=1 | 24.70 out tok/s, TPOT 37.90 ms, **8/8 ok** | +| c=2 | **VOID — ok=1, failed=7** | + +``` +engine-fatal: EngineCore busy loop threw: vt: propose_drafts_block: context position +discontinuity (accumulation out of sync with the target's committed positions) +at src/vllm/v1/worker/gpu/runner.cpp:2961 +``` + +and every later request on that server returns 500 `[request submitted to a +stopped AsyncLLM]`. Memory was not a factor; the box had ample headroom at c=2. + +The operator then ran the isolation. **With `--speculative-config` omitted and +everything else identical, the same two concurrent requests both complete** +(`req0: ok completion=32`, `req1: ok completion=32`, `fatal_lines=0`). So +batching, scheduling, the paged KV cache, the block tables and the sampler all +serve two sequences correctly. The defect is confined to the DFlash2 draft's +per-request context accumulation. + +**`ok=1` is the signature of the mechanism below**, not an incidental count. Two +requests are admitted together; the first completes; the second dies on the very +next step; the six that had not started yet then meet a stopped engine. + +## Why the accumulation desynchronises + +`GPUModelRunner` holds the draft context in four arrays indexed by **batch row** +(`include/vllm/v1/worker/gpu/runner.h:852-870`): + +```cpp +std::vector> dflash_kv_store_; +std::vector dflash_ctx_len_; +std::vector dflash_ctx_reqid_; +std::vector dflash_ctx_disabled_; +``` + +A row index is **not** stable for a request's lifetime in this tree. +`InputBatch::condense` slides a live request down into the hole a finished +neighbour left (`src/vllm/v1/worker/gpu/input_batch.cpp:611-760`; the row move is +`:686-706`), and `InputBatch::swap_states` exchanges two live rows +(`:762-847`). Both permute every per-slot array they know about — `req_ids`, +`num_computed_tokens_cpu`, `num_accepted_tokens`, `last_sampled_tokens`, +`prefill_len`, the block-table rows (`:706` `block_table.move_row`), the +index-keyed sampling maps. They know nothing about the four above, because those +live in `GPUModelRunner`, not in `InputBatch`. + +So when the first of two concurrent requests finishes, condense moves the +survivor from row 1 into row 0, and the survivor meets row 0's bookkeeping: + +1. `dflash_ctx_reqid_[0]` still names the departed request, so the reuse test at + `src/vllm/v1/worker/gpu/runner.cpp:2895-2906` reads a changed occupant. +2. It therefore allocates a **fresh empty store** and sets + `dflash_ctx_len_[0] = 0` — discarding a context the survivor is still using. +3. The survivor is mid-decode at absolute position L > 0, so the invariant at + `runner.cpp:2939-2945` sees `step.positions[rows[0]] == L_target != 0` and + refuses. + +**The invariant is correct and stays.** It is the only reason this is a loud +failure rather than a silent one: the alternative to the throw is drafting from a +context belonging to a different request. What is wrong is that the state it +guards is keyed by something the batch is free to change underneath it. + +This tree already names this exact bug class for a different array, in a comment +sitting above the log that was written to solve it +(`include/vllm/v1/worker/gpu/input_batch.h:240-252`): + +> Upstream needs no equivalent because it never condenses: `states.py:132` +> returns a finished request's slot to a free list and the slot index is stable +> for the request's lifetime. This log is the price of our condensed dense batch, +> not a deviation in what the state MEANS. + +`last_sampled_tokens` got a `LastSampledOp` log so its device mirror could follow +a row move. The DFlash2 arrays never got the equivalent treatment, and nothing +in the tree tied them to a request. + +### Why c=1 works and hides it + +At concurrency 1 the only row is row 0. A request finishes, row 0 empties, the +next request is admitted into row 0, the reuse test resets, and the reset is +**correct** — the new occupant is a fresh prefill whose first position is 0, so +the invariant passes. The condense move that breaks the state only exists when a +second live request has to be slid down over a departed one. + +### The `P == 1` capture gate is a consequence, not the cause + +`src/vllm/model_executor/models/qwen3_dflash.cpp:1577` admits the capture-safe +paged path only when `P == 1`; above that a fallback re-materialises each +request's context from the paged store every propose step. That is a +**performance** boundary. It is downstream of this defect — no batch ever reaches +`P == 2` for more than a step or two today, because the first row move kills the +engine. It is out of scope here and stays owed. + +## Upstream + +Read at `b389ac29465b33f9e9c534df221ea3c129e9793f` in `/home/mudler/_git/vllm`, +which is **beyond our parity pin `5559679229`** — beyond-pin vLLM, not a secondary +oracle. The clone is shallow and checked out at the pin, so every anchor below +was read with `git show b389ac2946:` and every line number is for the blob +at that revision, not for what is on disk. + +**Upstream has no host-side per-row draft context length at all.** The DFlash and +DFlash2 speculators are stateless across steps with respect to context: every +step re-derives the draft's whole KV addressing from the **target's own +`positions` array**, and writes draft KV into a paged KV cache group at the slot +addressed by that absolute position. + +- The draft's context lives in a paged KV cache group sharing the target + runner's `BlockTables` + (`vllm/v1/worker/gpu/spec_decode/dflash/speculator.py:181-193`, + `vllm/v1/worker/gpu/spec_decode/speculator.py:217-223`). Context KV is + re-inserted per step by `precompute_and_store_context_kv` + (`dflash/speculator.py:434-438`). +- Addressing is by absolute position, never by a running append cursor: + `ctx_pos = target_positions[ctx_start + j]` then + `ctx_block_num = ctx_pos // (block_size * CP_SIZE)` + (`dflash/speculator.py:562-590`). +- The one anchor is re-read from the target every step: + `last_valid_pos = tl.load(target_positions_ptr + valid_ctx_end - 1)` + (`dflash/speculator.py:553`), with the rejected suffix trimmed at `:542-544`. + The draft block's positions are `last_valid_pos + 1 + query_off` (`:593`). +- **Every tensor that outlives a step is indexed by the persistent request + slot**, resolved inside the kernel as + `req_state_idx = tl.load(idx_mapping_ptr + req_idx)` + (`dflash/speculator.py:536`) and carried forward through `sample_idx_mapping` + (`:639`). DFlash2's own cross-step state, `_cached_candidate_ids`, is written + at `cache_base = (req_state * num_steps + step) * top_k` + (`dflash2/speculator.py:95`), never at a row index. Its step-scoped + `_selector_scores` is row-indexed and written-then-read inside one step + (`dflash2/speculator.py:120-129`, `:213`, `:215`). +- The V2 runner those speculators live in **has no `condense` and no + `swap_states`**; a finished request's slot returns to a free list and the slot + is stable for the request's lifetime (`vllm/v1/worker/gpu/states.py:29,100,132`), + with the batch-row view produced by a per-step gather + (`vllm/v1/worker/gpu/block_table.py:143-170`). +- In the **legacy V1** runner, where rows *are* condensed, the draft's block-table + row is moved with the request: + `vllm/v1/worker/gpu_input_batch.py:786` `self.block_table.move_row(...)` fans + out to every KV group including the draft's + (`vllm/v1/worker/block_table.py:367-373`). `DFlashProposer` itself holds only + step-scoped buffers (`vllm/v1/spec_decode/dflash.py:49-72`, `:138`, `:142`, + `:292-299`). +- Upstream's own test already pins the row-to-slot indirection: + `tests/v1/spec_decode/test_dflash_prepare_inputs.py:46` passes + `idx_mapping=torch.tensor([2])` and `:139` asserts `sample_idx_mapping[:3] == + [2,2,2]` — a batch row deliberately mapped to a non-identity persistent slot. + +**Both upstream shapes key the draft context to the REQUEST.** V2 does it with a +stable slot plus an `idx_mapping` indirection; V1 does it by moving the draft's +block-table row along with the request. Ours does neither. That is the whole +defect. + +## Design + +**Key the four arrays by request id.** One `std::unordered_map` on the runner, holding the store, the context length and the +disabled flag; the row loop resolves a pointer per row once per step and every +existing use reads through it. + +This is upstream's V2 invariant expressed in our structure, not a workaround for +it. Upstream's persistent draft state is indexed by a key that survives any +reordering of the batch; a request id is such a key here, and it is the same key +`InputBatch::req_id_to_index` already uses. Every row permutation the batch can +perform — today's `condense` and `swap_states`, and any future one — becomes a +no-op for the draft context, which is the property that makes upstream's V2 +speculator indifferent to row order in the first place. + +```cpp +struct DflashReqCtx { + std::shared_ptr store; + int32_t ctx_len = 0; + bool disabled = false; +}; +std::unordered_map dflash_ctx_; +``` + +Three consequences, all of them simplifications: + +1. The reuse test at `runner.cpp:2895-2906` **disappears**. "Has this row's + occupant changed" was only ever a proxy for "is this state this request's"; + with the map the question cannot be asked wrongly. A first sight of a request + id constructs the entry, which is the reset. +2. `dflash_ctx_disabled_`'s own comment — "the flag is a property of the REQUEST, + not of the row" (`runner.cpp:2904-2905`) — becomes literally true instead of + approximately true. +3. The **decode-first reorder** is fixed for free, and this is a second live + trigger rather than a hypothetical one. `reorder_batch_to_split_decodes_and_prefills` + runs UNCONDITIONALLY on every step (`runner.cpp:1324`) and swaps live rows + through `swap_states` (`:207`) to put decode -> short_extend -> long_extend -> + prefill in that order. It is a no-op only while the batch's arrival order + already satisfies that ordering — which is the common case, because condense + keeps older decoding requests at low rows and a new arrival appends at the end + as a prefill, so `req_regions == target_regions` and no swap is emitted. That + is why the #2008 measurement met the condense move first and not this. A batch + whose regions are out of order — an older row still prefilling while a newer + one decodes — does emit swaps, and on the pre-change code those swapped two + live requests' draft contexts onto each other's rows. + + An earlier draft of this spec called that path "inert for a Qwen3 DFlash2 + target today". That was wrong, and it is corrected here rather than quietly: + the reorder has no model-family gate at its call site. + +The entries are pruned each step against `InputBatch`'s own membership, so a +finished or preempted request releases its device store on the step after it +leaves the batch. `InputBatch` is the authority on residency, not +`exec_state_.req_ids`, which lists only the rows scheduled this step. + +### What was rejected + +**Adopting upstream's paged shape** — registering the draft's context as a KV +cache group and letting `MultiGroupBlockTable::move_row` carry it, which is +exactly how upstream's legacy V1 path stays correct — is the right end state and +is not this change. It replaces `DflashDeviceKVStore` entirely, changes how the +draft's capacity is budgeted (#1919's sizing, and #2007's two-pool split), and +removes the private store the W8/W11 device-propose and paged-attention fast +paths are written against. It is a wave, not a bug fix, and it would land the +concurrency repair behind a rewrite. Recorded under `## Owed`. + +**Permuting the arrays from a `LastSampledOp`-style log** was the other candidate +and is worse on two counts. It keeps the row-indexed representation whose only +defect is that it is row-indexed, so every future batch mutation owes it another +entry; and the existing log is drained and cleared by +`replay_last_sampled_ops` (`runner.cpp:3456-3492`, cleared at `:1537` and +`:3431`), so a second consumer would race the first for the same buffer. + +## Risks + +- **A masking repair passes the obvious gate.** Marking the moved row disabled + (the #1919 fallback at `runner.cpp:2917-2921`) `continue`s *before* the + invariant, so it makes the throw go away, and the tokens are **identical** — + the verify is lossless, so a request that stops speculating emits exactly what + it emitted before, only slower. A token gate cannot see the difference. The + gate below therefore asserts that the survivor keeps proposing, and mutation A′ + measures that this refuses the repair rather than asserting that it would. +- **Per-step hashing on the propose path.** One string lookup per request per + step, bounded by `max_num_seqs`. At the measured TPOT of 37.90 ms this is + unmeasurable; it is stated so the ladder below can falsify it. +- **A leaked store is a device allocation.** Pruning is part of the change, not a + follow-up. + +## Tests + +`tests/vllm/v1/spec_decode/test_dflash2_concurrency.cpp`, its own binary for the +reason every fixture binary here has one: `VT_SPEC_TRACE` is latched once per +process by a function-local static on the first propose. + +Two concurrent requests are driven through the **synchronous production front** +(`LLMEngine::add_request` + `step()`), so the step boundaries the trace reports +are the ones the scheduler took. The neighbour is added first (row 0) and asks +for one token; the survivor is added second (row 1) and asks for ten. The +neighbour finishes first, condense slides the survivor into row 0, and the next +propose meets the defect. + +Four legs: + +1. **The engine does not throw.** Red-before this change with + `propose_drafts_block: context position discontinuity`. +2. **Concurrency was actually reached** — some step reports `rows=2`. Without + this the case would pass on a build that served the two requests one after the + other and never exercised a move. +3. **The survivor never stops proposing** — no `[spec-propose] NO proposing rows + this step` line (`runner.cpp:3259-3262`). This is the leg that refuses the + fallback masking repair. +4. **It proposes at every step it is alive for** — a count, not an absence. The + concurrent run's solo-row proposes are the solo run's minus the one step it + shared with the neighbour. + +Plus token-exactness of the survivor's output against the solo control, which is +necessary and not sufficient. + +### The leg that is deliberately absent, and why + +The obvious fifth leg is to compare the drafted BLOCKS either side of the move +against the solo run's, on the argument that a repair which resets the context +instead of moving it drafts from an empty context and is caught while every token +stays unchanged. **That leg was written, run, and removed: it is a tautology on +this fixture.** With the production invariant deleted and the context reset at +every row move, this draft still emits `12 12 12` at every step of both runs. The +synthetic draft's block is constant in the context — seeded-noise weights over a +24-token vocabulary, and the selector walk collapses to one id — so the +comparison asserts a constant against itself, and nine passing `CHECK`s measured +nothing. It is recorded here rather than silently dropped because the next person +to reach for that leg will find the same fixture. +(`test_dflash2_runner_reach`'s value-sensitivity case is unaffected: it moves the +drafts by changing the selector's WEIGHTS, not the context.) + +What pins the survivor's context to the survivor is therefore the **two +production invariants this change leaves alone**, standing with legs 1 and 3. No +throw means that for every proposing row `positions[rows[0]] == ctx_len` **and** +`ctx_len == DeviceKVNumCtx(store)`; still proposing means the row reached those +invariants rather than skipping them down the disabled path. Together they say +the survivor proposed with a context length equal to its own committed position, +held by a store containing exactly that many rows. This is why the row's method +depends on the invariants staying, and not only as a matter of policy. + +### Mutation evidence + +| Mutation, built on the PRE-CHANGE code | Result | +|---|---| +| **A′ — the fallback repair.** On a position mismatch, mark the row disabled and fall back instead of asserting. Makes the throw stop; emits identical tokens. | **CAUGHT.** Case 1 goes green, and leg 3 reads `none_lines == 9` — the survivor stopped proposing for all nine remaining steps. This is the token-invisible repair, refused by measurement. | +| **B′ — the "silence the check" repair.** Delete the position invariant and let the reset context stand. | **NOT CAUGHT** by this file, which is how the tautology in leg 5 was found. Recorded under `## Owed`: the invariant is load-bearing and nothing in this row's gate holds it. | +| **Reachability.** Delete the `propose_drafts_block` call site in `propose_drafts_dflash`. | See `## Gates`. | + +## Gates + +```sh +cmake -S . -B build -DVLLM_CPP_BUILD_TESTS=ON -DCMAKE_EXPORT_COMPILE_COMMANDS=ON +cmake --build build -j 4 +ctest --test-dir build --output-on-failure +``` + +Focused: `ctest --test-dir build -R 'dflash' --output-on-failure`. + +Reachability mutation (`.agents/reachability.md`): delete the production call +site and rerun the focused gate. The capability enters through +`LoadedEngine` -> `LLMEngine::add_request`/`step` -> `EngineCore::step` -> +`GPUModelRunner::execute_model` -> `propose_drafts_block`; no test constructs the +runner or the store by hand. + +### The end-to-end measurement, and what falsifies it + +The concurrency ladder at the #2008 flags: `--num-blocks 3744 --max-num-seqs 16 +--max-model-len 8192 --no-enable-prefix-caching --speculative-config +'{"method":"dflash","model":"/draft","num_speculative_tokens":8}'`, `vllm bench +serve --backend openai-chat`, in 1024 / out 512, rungs c = 1, 2, 4, 8, 16. +**The operator runs this. Predicted before the run:** + +| c | predicted out tok/s | predicted TPOT (ms) | vLLM | SGLang | +|---:|---:|---:|---:|---:| +| 1 | 24.7 (unchanged) | 37.9 (unchanged) | 24.36 | 25.20 | +| 2 | 38-46 | 40-48 | 38.59 | 44.93 | +| 4 | 60-78 | 45-58 | 64.25 | 77.06 | +| 8 | 75-105 | 60-85 | 80.95 | 109.24 | +| 16 | 85-135 | 90-150 | 99.87 | 142.61 | + +The bands are wide on purpose and the reason is stated rather than hidden: this +change makes the rungs *exist*, and what they measure once they do is the +`P == 1` capture gate at `qwen3_dflash.cpp:1577` and the two-pool allocation +of #2007, neither of which this row touches. A rung that lands at the bottom of +its band is that fallback path being measured for the first time, not this fix +underperforming. + +**This must NOT merge on any of these:** + +- **c=1 regresses at all** — 24.7 out tok/s or TPOT 37.9 ms moving outside noise + in the wrong direction. c=1 never takes a condense move with a live neighbour, + so this change must be inert there. A c=1 regression means the per-step map + lookup or the pruning scan is on a hotter path than claimed, and the claim is + wrong. +- **Any rung still VOIDs**, on this or any other invariant. A ladder that reaches + c=4 and dies at c=8 is a second defect this change did not find, and it stays + open rather than merging behind a partial result. +- **c=2 comes in below c=1.** Two sequences that each go slower than one sequence + alone would mean the draft is serialising the batch, which is a different + defect from the one diagnosed here. +- **A rung completes with `failed > 0`**, or any `engine-fatal` line, at any + concurrency. + +A rung that lands *below its band but above c=1, with `failed == 0` and no fatal +line*, is not a falsifier: it is #2007 and the `P == 1` gate being measured, and +those are named, owned and out of scope. The distinction is exactly which claim +each result contradicts, and this row claims only that concurrency **works**. + +## Owed + +- The draft context as a real KV cache group carried by + `MultiGroupBlockTable::move_row`, which is upstream's own shape on both its + paths. Tracked by the row; not attempted here. +- The `P == 1` capture gate at + `src/vllm/model_executor/models/qwen3_dflash.cpp:1577`. Above one proposing + row the paged capture-safe route is refused and a fallback re-materialises each + request's context every propose step. This is the first change that lets a + batch reach `P > 1` at all, so it is also the first that makes this cost + measurable. +- [#2007](https://github.com/mudler/vllm.cpp/issues/2007) — attention and + recurrent state allocated from two pools, which is why c=32 at k=8 is + unservable. Interacts with the ladder above and is deliberately not widened + into. +- [**#2009**](https://github.com/mudler/vllm.cpp/issues/2009) — **the position + invariant at `runner.cpp:2939-2945` is itself ungated.** + Deleting it (mutation B′ above) leaves this row's own gate green, and this row + does not close that. It is the guard the whole draft-context accumulation rests + on and the one the operator required be kept, so "nothing would notice if it + went" is a real gap. Gating it needs a fixture whose draft is sensitive to its + context, which this one is not — see the absent leg above. Its own issue. + +- **UNVERIFIED, and labelled so deliberately: a prefix-cache hit or a resumed + request may trip the same invariant at c=1.** The reasoning is that a request + admitted with `num_computed_tokens > 0` has no draft context for the tokens the + cache supplied, so its first propose would read `step.positions[rows[0]] > 0` + against `L == 0` and the same `VT_CHECK` would refuse. **It was not + reproduced.** A throwaway probe on this fixture forced + `EngineParams::enable_prefix_caching = true` (which + `ResolveEnablePrefixCaching` honours verbatim, + `src/vllm/entrypoints/model_loader.cpp:1075-1078`) and issued the same + 20-token prompt twice; both requests completed, nothing threw, and the engine's + own `prefix_cache_metrics()` reported **`queries=40 hits=0`**. The cache never + engaged, so the probe measured nothing about the hypothesis — the target here + is a GDN hybrid, the family upstream defaults prefix caching OFF for. This + stays a reasoned hypothesis and is NOT reported as a defect. Confirming it + needs a decoder-only DFlash2-capable target, which this fixture is not. + + Recorded because the answer, if it is real, is upstream's empty-draft path + (`ngram_proposer.py:156-159`) — which is also the shape of the masking branch + this row's gate exists to refuse, so it would have to be a separate change with + its own discriminating gate either way. + +## Now + +`SPEC-DFLASH2` stays `ACTIVE`. diff --git a/include/vllm/v1/worker/gpu/runner.h b/include/vllm/v1/worker/gpu/runner.h index 042bd5f2b..67ccd4426 100644 --- a/include/vllm/v1/worker/gpu/runner.h +++ b/include/vllm/v1/worker/gpu/runner.h @@ -854,41 +854,76 @@ class GPUModelRunner final : public ModelRunnerBase { bool dspark_sample_from_anchor_ = true; bool use_dspark() const { return dspark_weights_ != nullptr; } // Per-request PERSISTENT context KV store (D9 persistent paged draft-KV — the - // perf form of vLLM's incrementally-written draft KV cache). dflash_kv_store_[i] - // holds request i's per-layer bf16 context K/V (K normed+RoPE'd, V raw) for its - // committed positions 0..L_i-1 (L_i = dflash_ctx_len_[i]). Each verify step - // projects ONLY the newly-accepted rows (AppendContextKVHost) and APPENDS them, - // instead of re-projecting the whole growing context (the D5/D7 O(context^2) - // recompute). Bit-identical to the recompute by per-row projection independence. - // dflash_ctx_reqid_[i] tracks the occupant so a reused batch slot resets its - // store; rejected drafts' rows are never appended (rollback = don't-append). - // Indexed by the runner's condensed-dense batch row. Sized on set_dflash_draft. - // D11 A-wire: the store is now the DEVICE-RESIDENT append-only draft-KV store - // (DflashDeviceKVStore, opaque, one shared_ptr per condensed-dense batch row). - // AppendContextKVDevice keeps the projected bf16 K/V on-device (no D<->H round - // trip) and ForwardBlockLogitsWithDeviceKV runs the block forward straight off - // the device store — bit-identical to the D9 host path, and the capture-ready - // substrate for Parts B/C. shared_ptr-to-incomplete is safe: MakeDeviceKVStore - // constructs the control block (with its deleter) in qwen3_dflash.cpp. - std::vector> dflash_kv_store_; - std::vector dflash_ctx_len_; - std::vector dflash_ctx_reqid_; - // #1919: the store's resolved capacity, taken ONCE at set_dflash_draft from - // this engine's own max_model_len, and the per-row "this request no longer - // fits" flag. + // perf form of vLLM's incrementally-written draft KV cache). One entry holds + // that request's per-layer bf16 context K/V (K normed+RoPE'd, V raw) for its + // committed positions 0..ctx_len-1. Each verify step projects ONLY the + // newly-accepted rows (AppendContextKVDeviceRows) and APPENDS them, instead of + // re-projecting the whole growing context (the D5/D7 O(context^2) recompute). + // Bit-identical to the recompute by per-row projection independence; rejected + // drafts' rows are never appended (rollback = don't-append). + // + // D11 A-wire: the store is the DEVICE-RESIDENT append-only draft-KV store + // (DflashDeviceKVStore, opaque). AppendContextKVDevice keeps the projected bf16 + // K/V on-device (no D<->H round trip) and ForwardBlockLogitsWithDeviceKV runs + // the block forward straight off it — bit-identical to the D9 host path, and + // the capture-ready substrate for Parts B/C. shared_ptr-to-incomplete is safe: + // MakeDeviceKVStore constructs the control block (with its deleter) in + // qwen3_dflash.cpp. + // + // KEYED BY REQUEST ID, NOT BY BATCH ROW (#2008). These three fields were three + // arrays indexed by the runner's condensed-dense batch row, with a fourth + // recording each row's occupant so a reused slot could reset. A row index is + // not stable for a request's lifetime here: `InputBatch::condense` slides a + // live request down into the hole a finished neighbour left, and `swap_states` + // exchanges two live rows. Both permute every per-slot array they own — + // including the block-table rows — and neither knows these exist, because they + // live on the runner rather than in `InputBatch`. So the survivor of a + // completed pair met the departed request's bookkeeping, the occupant test + // read a changed id, the store was reset to EMPTY under a request still using + // it, and `propose_drafts_block`'s position invariant then refused. #2008 + // measured what that costs: DFlash2 served c=1 at 24.70 out tok/s and VOIDed + // at c=2 with ok=1, after which every later request on that server came back + // `[request submitted to a stopped AsyncLLM]`. // - // The flag is STICKY for the lifetime of the request occupying the row, and - // that is forced rather than chosen. `propose_drafts_block` keeps - // `dflash_ctx_len_` in lockstep with the store's `num_ctx` and asserts both - // against the target's committed positions; a step that declines to append - // breaks that lockstep, so every later step for the same request must decline - // too. It is cleared where the store is rebuilt — when a reused dense slot - // changes occupant — so a later, shorter request on the same row speculates - // normally. Upstream's own skip is monotone in the same way: its - // `num_tokens >= max_model_len` condition only ever becomes true - // (`vllm/v1/spec_decode/ngram_proposer.py:156-159`). + // Upstream keys the same state to the request on both of its paths. Its V2 + // runner, where DFlash2 lives, has no `condense` at all — a finished request's + // slot returns to a free list and stays that request's for its lifetime + // (`vllm/v1/worker/gpu/states.py:29,100,132` @ `b389ac2946`) — and every + // cross-step speculator tensor is indexed through + // `req_state_idx = idx_mapping[req_idx]` + // (`vllm/v1/worker/gpu/spec_decode/dflash/speculator.py:536`). Its legacy V1 + // runner does condense, and there the draft's block-table row moves with the + // request (`vllm/v1/worker/gpu_input_batch.py:786` -> + // `vllm/v1/worker/block_table.py:367-373`). A request id is the key that + // survives any reordering of OUR batch, so every permutation the batch can + // perform is a no-op here — which is the property that makes upstream's + // speculator indifferent to row order in the first place. + // + // Entries are pruned each propose against `InputBatch`'s own membership, so a + // finished or preempted request releases its device store on the step after it + // leaves the batch. + struct DflashReqCtx { + std::shared_ptr store; + // Committed context length L. Kept in lockstep with the store's own num_ctx, + // and asserted against it every propose (SPEC-DFLASH2 W8, #1838). + int32_t ctx_len = 0; + // #1919: this request no longer fits the store and runs on the target alone. + // STICKY for the request's lifetime, and that is forced rather than chosen: + // `propose_drafts_block` keeps `ctx_len` in lockstep with the store's + // `num_ctx` and asserts both against the target's committed positions, so a + // step that declines to append breaks that lockstep and every later step for + // the same request must decline too. Upstream's own skip is monotone in the + // same way: its `num_tokens >= max_model_len` condition only ever becomes + // true (`vllm/v1/spec_decode/ngram_proposer.py:156-159`). Being a property + // of the REQUEST — which the row-indexed form could only approximate, and + // its comment already claimed — it now simply ends with the request. + bool disabled = false; + }; + std::unordered_map dflash_ctx_; + // #1919: the store's resolved capacity, taken ONCE at set_dflash_draft from + // this engine's own max_model_len. The "no longer fits" flag it pairs with is + // `DflashReqCtx::disabled` above. vllm::Qwen3DFlashModel::DflashCtxStoreSizing dflash_ctx_sizing_; - std::vector dflash_ctx_disabled_; // Draft KV cache (`fa_draft` group) backing storage, owned by the runner and // allocated in initialize_kv_cache when spec is on. draft_attn_kv_ (declared // above) views into these buffers. Empty on the default path. diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp index 6d7d27de2..f7ed973af 100644 --- a/src/vllm/v1/worker/gpu/runner.cpp +++ b/src/vllm/v1/worker/gpu/runner.cpp @@ -2855,10 +2855,7 @@ void GPUModelRunner::set_dflash_draft(const vllm::Qwen3DFlashWeights* weights, dflash_tap_layer_ids_.push_back(id.get()); } } - dflash_kv_store_.clear(); - dflash_ctx_len_.clear(); - dflash_ctx_reqid_.clear(); - dflash_ctx_disabled_.clear(); + dflash_ctx_.clear(); // #1919: resolve the draft context store's capacity from THIS engine's // advertised context, and say so. Before this, the capacity was a @@ -2957,11 +2954,38 @@ void GPUModelRunner::propose_drafts_block( const int num_mask_rows = num_query_per_req - 1; const StepInputs& step = exec_state_.step; - if (static_cast(dflash_ctx_len_.size()) < num_reqs) { - dflash_kv_store_.resize(static_cast(num_reqs)); - dflash_ctx_len_.resize(static_cast(num_reqs), 0); - dflash_ctx_reqid_.resize(static_cast(num_reqs)); - dflash_ctx_disabled_.resize(static_cast(num_reqs), false); + // #2008 — THE DRAFT CONTEXT BELONGS TO THE REQUEST, NOT TO THE BATCH ROW. + // + // Release the entries of requests that have left the batch, then resolve one + // pointer per row for the rest of this function. `InputBatch` is the authority + // on residency, not `exec_state_.req_ids`, which lists only the rows SCHEDULED + // this step; a resident request that happens not to be scheduled must keep its + // context. The scan is bounded by `max_num_seqs` and runs once per step, off + // the per-token path. + // + // The pruning is not housekeeping: each entry owns a device allocation sized + // from `dflash_ctx_sizing_.slots`, so a leaked entry is leaked device memory. + if (dflash_ctx_.size() > static_cast(input_batch_.num_reqs())) { + for (auto it = dflash_ctx_.begin(); it != dflash_ctx_.end();) { + it = (input_batch_.req_id_to_index.count(it->first) == 0) + ? dflash_ctx_.erase(it) + : std::next(it); + } + } + // Resolving here replaces the reused-slot test this loop used to open with. + // "Has this row's occupant changed" was only ever a proxy for "is this state + // this request's", and the proxy is what #2008 broke: after a condense move + // the answer was yes for a request whose context was perfectly valid, and the + // reset threw it away. Keyed by id the question cannot be asked wrongly — a + // first sight of a request id constructs its entry, and THAT is the reset. + std::vector row_ctx(static_cast(num_reqs), nullptr); + for (int i = 0; i < num_reqs; ++i) { + DflashReqCtx& c = dflash_ctx_[exec_state_.req_ids[static_cast(i)]]; + if (c.store == nullptr) { + c.store = Qwen3DFlashModel::MakeDeviceKVStore(config, queue_, + dflash_ctx_sizing_.slots); + } + row_ctx[static_cast(i)] = &c; } // SPEC-DFLASH2 W8 (#1838): the propose pre-phase timer. Before W8 everything @@ -2998,29 +3022,21 @@ void GPUModelRunner::propose_drafts_block( std::vector fell_back(static_cast(num_reqs), false); for (int i = 0; i < num_reqs; ++i) { - // Reset a reused dense slot (a new request now occupies this row). - if (dflash_ctx_reqid_[static_cast(i)] != - exec_state_.req_ids[static_cast(i)]) { - dflash_kv_store_[static_cast(i)] = Qwen3DFlashModel::MakeDeviceKVStore( - config, queue_, dflash_ctx_sizing_.slots); - dflash_ctx_len_[static_cast(i)] = 0; - dflash_ctx_reqid_[static_cast(i)] = - exec_state_.req_ids[static_cast(i)]; - // A new occupant starts speculating again (#1919). The flag is a property - // of the REQUEST, not of the row. - dflash_ctx_disabled_[static_cast(i)] = false; - } + // #2008: no reused-slot reset here any more. This request's state was + // resolved by id above, so it is this request's whether or not the batch + // moved it, and a request seen for the first time started empty. + DflashReqCtx& ctx = *row_ctx[static_cast(i)]; // #1919: a request whose context has outgrown the store stops speculating // for the rest of its life, and this test comes BEFORE the two invariants // below because a disabled row stops maintaining both: it neither appends - // nor advances `dflash_ctx_len_`, so its counter and the target's committed + // nor advances its `ctx_len`, so that counter and the target's committed // positions legitimately diverge from here on. // // What such a row PROPOSES is decided in section 4, which is also where the // two scheduling modes part company; `fell_back` is how this loop tells it // which rows are in that state and are decoding rather than still // prefilling. - if (dflash_ctx_disabled_[static_cast(i)]) { + if (ctx.disabled) { fell_back[static_cast(i)] = !(i < static_cast(exec_state_.discard.size()) && exec_state_.discard[static_cast(i)]); @@ -3045,18 +3061,18 @@ void GPUModelRunner::propose_drafts_block( // Invariant: this step's first committed token sits at absolute position L // (== current context length). A violation means the accumulation lost sync // (the I5e async-input-combine bug class) — assert rather than corrupt. - const int64_t L = dflash_ctx_len_[static_cast(i)]; + const int64_t L = ctx.ctx_len; VT_CHECK(step.positions[static_cast(rows[0])] == L, "propose_drafts_block: context position discontinuity (accumulation " "out of sync with the target's committed positions)"); // SPEC-DFLASH2 W8 (#1838): the runner's counter and the DEVICE store must // agree, or the append is dead. The W8 mutation run proved the check above - // cannot see that state: with the append call deleted, `dflash_ctx_len_` + // cannot see that state: with the append call deleted, `ctx_len` // kept advancing, the store stayed empty, every propose ran CONTEXT-FREE, // and every gate stayed green — well-formed drafts, lossless verify, only // ACCEPTANCE falls, the exact invisible-defect class this row exists to // remove. This host integer comparison is what makes that state loud. - VT_CHECK(L == Qwen3DFlashModel::DeviceKVNumCtx(*dflash_kv_store_[static_cast(i)]), + VT_CHECK(L == Qwen3DFlashModel::DeviceKVNumCtx(*ctx.store), "propose_drafts_block: the runner's context length and the device " "store's num_ctx disagree — the context-KV append is dead or " "double-run (SPEC-DFLASH2 W8, #1838)"); @@ -3087,10 +3103,9 @@ void GPUModelRunner::propose_drafts_block( // worth of context and keeps the W11 paged fast route, which // `ClassifyDflashBlockAttn` would otherwise drop below its capacity // conjunct. - const int64_t capacity = Qwen3DFlashModel::DeviceKVCapacity( - *dflash_kv_store_[static_cast(i)]); + const int64_t capacity = Qwen3DFlashModel::DeviceKVCapacity(*ctx.store); if (L + append + num_query_per_req > capacity) { - dflash_ctx_disabled_[static_cast(i)] = true; + ctx.disabled = true; std::cerr << "vllm.cpp: request " << exec_state_.req_ids[static_cast(i)] << " has outgrown the draft speculative context (" << (L + append) << " context tokens plus a " << num_query_per_req @@ -3111,10 +3126,9 @@ void GPUModelRunner::propose_drafts_block( new_rows[static_cast(j)] = static_cast(rows[static_cast(j)]); new_pos[static_cast(j)] = static_cast(L + j); } - Qwen3DFlashModel::AppendContextKVDeviceRows(*dflash_kv_store_[static_cast(i)], - combined.tensor, new_rows, new_pos, backbone, - config, queue_); - dflash_ctx_len_[static_cast(i)] = static_cast(L + append); + Qwen3DFlashModel::AppendContextKVDeviceRows(*ctx.store, combined.tensor, new_rows, + new_pos, backbone, config, queue_); + ctx.ctx_len = static_cast(L + append); // A discarded (still-prefilling chunk) row commits its chunk's features but // proposes no draft — it has no valid last_sampled anchor yet. @@ -3124,7 +3138,7 @@ void GPUModelRunner::propose_drafts_block( // Block: anchor = last_sampled (the bonus/last committed token, re-embedded), // then k mask tokens; positions L' .. L'+k (L' = the new context length). - const int64_t Lp = dflash_ctx_len_[static_cast(i)]; + const int64_t Lp = ctx.ctx_len; const int32_t anchor = input_batch_.last_sampled_tokens[static_cast(i)]; blk_ids.push_back(anchor); blk_pos.push_back(static_cast(Lp)); @@ -3251,7 +3265,7 @@ void GPUModelRunner::propose_drafts_block( int64_t total_ctx = 0; for (int r = 0; r < P; ++r) { vllm::DflashDeviceKVStore* st = - dflash_kv_store_[static_cast(propose_rows[static_cast(r)])].get(); + row_ctx[static_cast(propose_rows[static_cast(r)])]->store.get(); stores.push_back(st); total_ctx += Qwen3DFlashModel::DeviceKVNumCtx(*st); ctx_cu.push_back(static_cast(total_ctx)); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 21d878869..e8bdf3c95 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1678,6 +1678,13 @@ target_include_directories(test_dflash2_embed_dedup_reach PRIVATE ${CMAKE_SOURCE # above 4096, which no other binary drives. vllm_cpp_add_test(test_dflash2_ctx_capacity vllm/v1/spec_decode/test_dflash2_ctx_capacity.cpp) +# SPEC-DFLASH2 (#2008): the draft context must follow the REQUEST, not the batch +# ROW. Its OWN binary for the reason every fixture binary here has one -- +# `VT_SPEC_TRACE` is latched once per process by a function-local static on the +# first propose, and this one needs level 1 with `max_num_seqs` above 1, which no +# other binary drives. +vllm_cpp_add_test(test_dflash2_concurrency + vllm/v1/spec_decode/test_dflash2_concurrency.cpp) target_include_directories(test_dflash_propose PRIVATE ${CMAKE_SOURCE_DIR}/src) # SPEC-NGRAM (ROAD-V1-D3) — the draft-free n-gram matcher unit gate (ports # vllm/tests/v1/spec_decode/test_ngram.py, host-side, runs everywhere). diff --git a/tests/vllm/v1/spec_decode/test_dflash2_concurrency.cpp b/tests/vllm/v1/spec_decode/test_dflash2_concurrency.cpp new file mode 100644 index 000000000..fd11c4353 --- /dev/null +++ b/tests/vllm/v1/spec_decode/test_dflash2_concurrency.cpp @@ -0,0 +1,233 @@ +// SPEC-DFLASH2 ([#2008](https://github.com/mudler/vllm.cpp/issues/2008)) — the +// DFlash2 draft context must belong to the REQUEST, not to the batch ROW. +// +// WHAT #2008 MEASURED. A DFlash2 K=8 server on an idle leased GB10 answers +// concurrency 1 at 24.70 out tok/s (8/8 ok) and DIES at concurrency 2, one +// request in: +// +// vt: propose_drafts_block: context position discontinuity (accumulation +// out of sync with the target's committed positions) +// +// after which every later request on that server returns +// `[request submitted to a stopped AsyncLLM]`. The operator then ran the +// isolation: with `--speculative-config` OMITTED and everything else identical, +// the same two concurrent requests both complete. So batching, scheduling, the +// paged KV cache, the block tables and the sampler all serve two sequences +// correctly, and the defect is confined to the draft's per-request context +// accumulation. +// +// WHY IT DESYNCHRONISES. `GPUModelRunner` holds the draft context in four +// arrays indexed by BATCH ROW — `dflash_kv_store_`, `dflash_ctx_len_`, +// `dflash_ctx_reqid_` and `dflash_ctx_disabled_`. A row index is not stable for +// a request's lifetime here: `InputBatch::condense` slides a live request down +// into the hole a finished neighbour left (input_batch.cpp:611-760, the row move +// at :686-706), and `InputBatch::swap_states` exchanges two live rows +// (input_batch.cpp:762-847). Both permute every per-slot array they know about +// — and they know nothing about the runner's four, which live in +// `GPUModelRunner`, not in `InputBatch`. +// +// So when the first of two concurrent requests finishes, condense moves the +// SURVIVOR from row 1 to row 0, and the survivor now meets row 0's bookkeeping: +// `dflash_ctx_reqid_[0]` still names the departed request, the reuse test at +// runner.cpp:2895-2906 reads a changed occupant, resets the store and sets +// `dflash_ctx_len_[0] = 0` — and the survivor is mid-decode at absolute +// position L > 0. The invariant at runner.cpp:2939-2945 then refuses, which is +// exactly what it is for. The header comment two lines above the arrays already +// names the shape of this bug for a DIFFERENT array +// (input_batch.h:240-252): "Upstream needs no equivalent because it never +// condenses: states.py:132 returns a finished request's slot to a free list and +// the slot index is stable for the request's lifetime." The DFlash2 arrays never +// got that treatment. +// +// UPSTREAM. vLLM's DFlash draft keeps no runner-side per-row context length at +// all: its context K/V lives in the engine's own paged KV cache +// (`vllm/model_executor/models/qwen3_dflash.py`), addressed through the block +// table, whose rows condense DOES move (`block_table.move_row`, +// input_batch.cpp:706). The context is per-REQUEST upstream. Keying our four +// arrays by request id is that same semantics expressed in our structure. +// +// WHAT THIS FILE GATES, AND WHY "IT DID NOT THROW" IS NOT ENOUGH. There is a +// second way to make the throw stop: mark the moved row disabled and let the +// request run on the target alone, which is the #1919 fallback and which +// `continue`s BEFORE the invariant. That produces IDENTICAL TOKENS — the verify +// is lossless, so a request that stops speculating emits exactly what it emitted +// before, only slower. A token gate cannot see it. So the cases below assert +// that the survivor KEEPS PROPOSING across the move, and at every step it is +// alive for — read off the production `VT_SPEC_TRACE` line, which prints the +// proposing-row count and prints `NO proposing rows this step` when that count +// is zero (runner.cpp:3259-3262). That is the leg which refuses the fallback +// repair, and it is MEASURED to refuse it rather than argued to: built on the +// pre-change code, that repair leaves case 1 green and turns this count into 9. +// +// The block-comparison leg a reader would expect next is deliberately absent, +// and the tail of the second case records why it would gate nothing here. +#include + +#include "dflash2_runner_fixture.h" + +#include "vllm/v1/engine/llm_engine.h" + +namespace { + +// Latched by a function-local static on the FIRST propose in the process, so it +// has to be set before any case runs. Level 1 is what prints `[spec-propose]`. +const bool kSpecTraceEnabled = [] { + ::setenv("VT_SPEC_TRACE", "1", 1); + return true; +}(); + +// One parsed `[spec-propose]` line. +struct ProposeLine { + int rows = 0; // P, the number of rows that proposed this step + std::string first; // the first proposing row's drafted ids + bool none = false; // the `NO proposing rows this step` variant +}; + +std::vector ParseProposeLines(const std::string& captured) { + std::vector out; + size_t at = 0; + const std::string tag = "[spec-propose] "; + while ((at = captured.find(tag, at)) != std::string::npos) { + const size_t body = at + tag.size(); + const size_t eol = captured.find('\n', body); + const std::string line = + captured.substr(body, (eol == std::string::npos) ? std::string::npos : eol - body); + ProposeLine p; + if (line.rfind("NO proposing rows", 0) == 0) { + p.none = true; + } else { + const size_t r = line.find("rows="); + if (r != std::string::npos) p.rows = std::atoi(line.c_str() + r + 5); + const size_t f = line.find("first=["); + if (f != std::string::npos) { + const size_t open = f + 7; + const size_t close = line.find(']', open); + if (close != std::string::npos) p.first = line.substr(open, close - open); + } + } + out.push_back(p); + at = (eol == std::string::npos) ? captured.size() : eol; + } + return out; +} + +// The survivor: a two-token prompt and enough output that it is still decoding +// long after its neighbour has left the batch. +constexpr int kKeeperTokens = 10; +// The neighbour: added FIRST, so it takes row 0 and the survivor takes row 1, +// and finishes FIRST, so condense slides the survivor down into row 0. With +// k = kSpecTokens = 3 a step can commit up to 4 tokens, so 1 is one step. +constexpr int kLeaverTokens = 1; + +struct RunResult { + std::string threw; + std::string keeper_text; + std::vector lines; +}; + +// Drive the SYNCHRONOUS production front (LLMEngine::add_request + step), so the +// step boundaries the trace reports are the ones the scheduler took and nothing +// depends on a drain deadline. +RunResult Run(bool with_neighbour) { + const HfConfig target = MakeDenseConfig(); + const ScratchDraftDir dir; + RunResult r; + const std::string captured = CaptureStderr([&] { + LoadedEngine eng(target, MakeDenseWeights(target), BuildFixture(), + DflashSpecParams(dir, /*max_model_len=*/0, /*max_num_seqs=*/4), + MakeDflash2Draft(target, /*muse_glimmer_scalars=*/false)); + vllm::v1::LLMEngine& e = eng.engine(); + try { + if (with_neighbour) e.add_request("leaver", "hello", Greedy(kLeaverTokens)); + e.add_request("keeper", "hello world", Greedy(kKeeperTokens)); + while (e.has_unfinished_requests()) { + for (const vllm::RequestOutput& o : e.step()) { + if (o.request_id == "keeper" && o.finished && !o.outputs.empty()) + r.keeper_text = o.outputs[0].text; + } + } + } catch (const std::exception& ex) { + r.threw = ex.what(); + } + }); + r.lines = ParseProposeLines(captured); + return r; +} + +} // namespace + +TEST_CASE("dflash2 #2008: a second request must not desynchronise the draft context") { + const RunResult conc = Run(/*with_neighbour=*/true); + INFO("threw: ", conc.threw); + CHECK(conc.threw.empty()); +} + +TEST_CASE("dflash2 #2008: the survivor keeps its context across the row move") { + const RunResult solo = Run(/*with_neighbour=*/false); + const RunResult conc = Run(/*with_neighbour=*/true); + INFO("solo threw: ", solo.threw); + INFO("conc threw: ", conc.threw); + REQUIRE(solo.threw.empty()); + REQUIRE(conc.threw.empty()); + + // The run genuinely reached concurrency: some step had TWO proposing rows. + // Without this the case would pass on a build that served the two requests + // one after the other and never exercised the move at all. + int max_rows = 0; + for (const ProposeLine& p : conc.lines) max_rows = std::max(max_rows, p.rows); + INFO("max proposing rows observed: ", max_rows); + CHECK(max_rows == 2); + + // The survivor never stops proposing. The #1919 fallback (mark the moved row + // disabled, run on the target alone) makes the throw go away and emits the + // SAME TOKENS, and this is the leg that separates it from a real fix: a step + // in which the only live request proposes nothing prints the `NO proposing + // rows` line. + int none_lines = 0; + for (const ProposeLine& p : conc.lines) none_lines += p.none ? 1 : 0; + CHECK(none_lines == 0); + + // And it proposes at EVERY step it is alive for, which is a count rather than + // an absence. The concurrent run's steps are the solo run's plus the one it + // shared with the neighbour, so the survivor's solo-row steps are one fewer. + // A repair that let the survivor skip a single propose — the shape a + // "re-seat it next step" patch takes — moves this number without moving a + // token. + std::vector conc_tail; + for (const ProposeLine& p : conc.lines) + if (!p.none && p.rows == 1) conc_tail.push_back(p.first); + std::vector solo_all; + for (const ProposeLine& p : solo.lines) + if (!p.none) solo_all.push_back(p.first); + REQUIRE_FALSE(conc_tail.empty()); + INFO("conc solo-row proposes: ", conc_tail.size(), " solo proposes: ", solo_all.size()); + CHECK(conc_tail.size() + 1 == solo_all.size()); + + // Token-exactness across the move. NECESSARY AND NOT SUFFICIENT, and this + // comment says which because the first version of this file got it wrong. + // + // WHAT IS **NOT** GATED HERE, MEASURED RATHER THAN ASSUMED. The obvious + // stronger leg is to compare the drafted BLOCKS either side of the move + // against the solo run's, on the argument that a repair which resets the + // context instead of moving it drafts from an empty context and is caught + // while every token stays unchanged. That leg was written, and it is a + // TAUTOLOGY on this fixture: with the production invariant deleted and the + // context reset on every row move, this draft still emits `12 12 12` at every + // single step of both runs. The synthetic draft's block is CONSTANT in the + // context — its weights are seeded noise over a 24-token vocabulary, and the + // selector walk collapses to one id — so a block comparison here asserts a + // constant against itself. (`test_dflash2_runner_reach`'s value-sensitivity + // case is not affected: it moves the drafts by changing the selector's + // WEIGHTS, not the context.) + // + // What pins the survivor's context to the survivor is therefore the two + // PRODUCTION invariants this change deliberately leaves alone, standing + // together with the two legs above. No throw (leg 1) means that for every + // proposing row `positions[rows[0]] == ctx_len` AND + // `ctx_len == DeviceKVNumCtx(store)`; still proposing (leg 3) means the row + // reached those invariants rather than skipping them down the disabled path. + // Together they say the survivor proposed with a context length equal to its + // OWN committed position, held by a store containing exactly that many rows. + CHECK(conc.keeper_text == solo.keeper_text); + CHECK_FALSE(conc.keeper_text.empty()); +}