diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index 34ad8fc2b6..6cf6a43eeb 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -174,7 +174,7 @@ it does not mean the global list contains that numeric target. | `BACKEND-CUDA-SM101` | CUDA compute 10.1 | `>=12.8,<13` globally; family via `f` on 13 | `CMakeLists.txt:105-118`; loose intersection `cmake/utils.cmake:393-481` | configurable target [CMakeLists.txt:37](../CMakeLists.txt#L37), applied globally at [CMakeLists.txt:64](../CMakeLists.txt#L64) | - | [CUDA inventory](specs/cuda-architecture-inventory.md); [arch additivity §W10](specs/cuda-arch-additivity.md) | `INVENTORIED` - **SCOPED (toolkit-blocked here), not build-supported.** nvcc 13.0 rejects `sm_101a` (`nvcc fatal: Unsupported gpu architecture 'sm_101a'`, measured); vLLM lists `10.1` only on its `<13` toolkit branch. Nothing arch-specific blocks it once a toolkit that accepts it is used (it would join the major-10 portable-only fan-out); build-verify deferred to that toolkit | - | | `BACKEND-CUDA-SM103` | CUDA compute 10.3 / `10.3a` (datacenter Blackwell) | `>=12.8,<13` globally; family via `f` on 13 | `CMakeLists.txt:105-118`; loose intersection `cmake/utils.cmake:393-481`. Fast paths NOT ported | **SINGLE-ARCH PORTABLE-KERNELS-ONLY (spec §W10):** all five fast-path cells resolve EMPTY; portable kernels only. Shares the identical major-10 portable bodies with the compiled `sm_100a` representative; no new code | **BUILD evidence (dgx, nvcc 13.0, `~/work/archexp`):** `103a` configures all five `DISABLED for [103a]`; the portable path is the SAME bodies compiled clean for the `sm_100a` representative; [CudaArchFeaturesTest.cmake](../cmake/CudaArchFeaturesTest.cmake) pins `103a`→all-features-EMPTY. **NO sm_103 execution here** | [CUDA inventory](specs/cuda-architecture-inventory.md); [arch additivity §W10](specs/cuda-arch-additivity.md); [datacenter fast-path](specs/cuda-arch-datacenter-fastpath.md) | `SPIKE` - **portable-only build-supported RETAINED (shares the major-10 portable bodies with the `sm_100a` representative); the datacenter-Blackwell FAST-PATH bodies are SPIKED for DERIVE-AND-SHIP per [datacenter fast-path](specs/cuda-arch-datacenter-fastpath.md)** (same tcgen05 CUTLASS legs as `sm_100a`, `10.3a` gencode). NOT runtime support; no board here | `CLAIM-CUDA-DATACENTER-SCOPE` | | `BACKEND-CUDA-SM110` | CUDA compute 11.0 (Blackwell family) | `>=13` | `CMakeLists.txt:105-118`; SM10/11 family rules at `:806-837,892-914`. Fast paths NOT ported EXCEPT `marlin-nvfp4`, ported 2026-08-11 | **SINGLE-ARCH PORTABLE-KERNELS + `marlin-nvfp4`, DATA-declared (spec §W10):** FOUR of the five fast-path cells resolve EMPTY; `marlin-nvfp4` resolves `ENABLED for [110]` since 2026-08-11 (the cell gained `11.0`; NO kernel source changed). Was "all five EMPTY" until then. No new code — the vendored dense/MoE Marlin bodies are the same ones sm_12x builds | **BUILD evidence (dgx, nvcc 13.0, `~/work/archexp`, `-DVLLM_CPP_TRITON=OFF`) — the compiled per-major-11 REPRESENTATIVE:** single-arch `110` configures all five `DISABLED for [110]`, compiles/links `-Werror` **0 warnings** (`EXIT=0`), `cuobjdump -lelf libvllm.a` shows **16 TUs carrying real `sm_110` cubins and nothing else**. Resolution asserted: [CudaArchFeaturesTest.cmake](../cmake/CudaArchFeaturesTest.cmake) pins `110`→all-features-EMPTY. **RUNTIME-VERIFIED on real sm_110 silicon 2026-07-27 (`CLAIM-CUDA-SM110-RUNTIME`, the FIRST non-GB10 runtime proof) — NVIDIA Jetson Thor (`ssh 192.168.68.23`, hostname `thor`, aarch64, JetPack R38, driver 580.00, `nvidia-smi compute_cap=11.0` CONFIRMING the inferred sm_110, nvcc `/usr/local/cuda-13.0` V13.0.48, cutlass ABSENT — not needed, all fast paths EMPTY).** Native build on Thor (`git archive` of `0f07fe34` → Thor, NOT rsync; `-DVLLM_CPP_CUDA_ARCHITECTURES=110 -DVLLM_CPP_CUDA=ON -DVLLM_CPP_TRITON=OFF`, Release): all five features `DISABLED for [110]`, compiles/links **0 warnings** under Release `-Werror` (CXX `-Wall -Wextra -Werror`; CUDA `--generate-code=arch=compute_110,code=[compute_110,sm_110]` + `-Werror=all-warnings`), `cuobjdump -lelf libvllm.a` shows **16 TUs of real `sm_110` SASS and nothing else** (the 22 fast-path TUs absent — the documented per-major-11 shape). **The portable bf16 forward RAN and is CORRECT:** `test_llama_paged_engine` (unsloth/Llama-3.2-1B bf16 safetensors, transferred dgx→Thor) drove the full paged LLMEngine on sm_110 and is **STRICT token-exact 12/16 prompts (192/192 tokens) vs the committed dgx-captured vLLM oracle greedy golden** — every prompt where vLLM's greedy is DETERMINISTIC — plus **15/16 bit-identical to the GB10 sm_121a anchor**; the remaining 4/16 are the ratified bf16 near-tie prompts (where GB10 itself already diverges from vLLM greedy; committed teacher-forced gap = 0.000 nats, exact ties), 3/4 of them bit-identical to the GB10 anchor. PASS under the [near-tie distributional gate](specs/cuda-arch-additivity.md). (The committed gate's hard anchor-REQUIRE encodes a GB10 bit-identity assumption predating any second board, so it aborts on p0's cross-arch near-tie tail branch; the underlying forward is correct.) **`marlin-nvfp4` KERNEL-RUNTIME-VERIFIED (numerics-gated) on real sm_110 silicon 2026-08-11 — the FEATURE-TABLE cell gained `11.0`; NO kernel source changed.** All 18 `__CUDA_ARCH__` guards under `src/vt/cuda/marlin/` enumerated: for `1100` each selects the same side an already-supported arch selects. Explicitly, sm_110 does **NOT** satisfy `__CUDA_ARCH__ < 890` (takes the `>= 890` side, as sm_90/sm_100/sm_12x do); both `< 890` sites are bare fp8-activation bail-outs with no `#else`, and this path pins `a_type = kBFloat16` so the predicate is false at compile time on every arch. **BUILD:** single-arch `110` configures `marlin-nvfp4: ENABLED for [110]`, 403/403 `EXIT=0`, including `sm80_kernel_bfloat16_fe2m1f_bfloat16.cu` (dense + MoE). **RAN** on Thor across 5 decode shapes x M in {1,2,4,8}, no illegal-instruction/misaligned faults. **NUMERICS vs the incumbent `vt::MatmulNvfp4`,** same operands, **random** fp4 weights + **random** per-group fp8-e4m3 scales + **random** bf16 activations (all three must be random: a constant activation makes any K-axis permutation produce the same dot product, and constant scales hide a scale-permutation bug — an earlier constant-input check of ours was vacuous and is retracted): 20 configurations / **631,424 elements**, **99.84%-100% of elements bit-for-bit equal**, remainder within one to a few bf16 ULP (largest abs deviation 4.0 on down_proj where the reference max magnitude is 1336 and the bf16 ULP there is 8, i.e. half an ULP), **0 elements above 5% relative error**, corr `1.00000000`. **NOT bitwise identical and not claimed as such** (only kv_proj at M=1/M=2 is 100%). **NEGATIVE CONTROL:** feeding Marlin the raw unpermuted `[N,K/16]` scales — the silent-repack-mismatch mode — is detected on **every element** (`corr=0.00000000`, `rel>5%: 8192/8192`), so the passing comparison is meaningful. **PERF:** 8.0x-29.0x per GEMM at M=1 on Qwen3-32B shapes (q 0.9461->0.0829 ms, kv 0.5009->0.0173, o 0.8093->0.0794, gate 2.3675->0.2943, down 2.6945->0.2870), and flat in M where the incumbent is linear. **END-TO-END SERVING (Qwen3-1.7B-NVFP4A16, two binaries differing ONLY in the CMake cell):** 16.61->81.63 tok/s at c=1 and 33.30->477.07 at c=8; concurrency scaling c=1->8 goes 2.00x -> 5.84x, i.e. the flat-with-concurrency curve is fixed. **SCOPE — what is NOT claimed:** this is throughput, **NOT a token-exact model-level gate** — no vLLM-oracle golden was run against the Marlin build, so this does **not** meet the `CLAIM-CUDA-SM110-RUNTIME` bar and does not upgrade it. No 32B end-to-end profile exists (a 32B attempt disrupted co-resident CUDA processes on the shared board and was abandoned; the default `--num-blocks 0` auto-sizing allocates ~27 GB regardless of model size). `8.0+PTX` stays UNCLAIMED — sm_8x has not been built or run. No dispatch policy changed: `kTileMinRows` untouched; the cell only makes the kernel available, and `MatmulNvfp4W4A16D` already routes to Marlin when `VT_MARLIN_NVFP4` is compiled in. | [CUDA inventory](specs/cuda-architecture-inventory.md); [arch additivity §W10](specs/cuda-arch-additivity.md) | `ACTIVE` - **RUNTIME-VERIFIED (portable bf16 path) on Thor sm_110 — the FIRST non-GB10 runtime proof.** 3-state SIGNAL: **(1) portable bf16/safetensors path = RUNTIME-VERIFIED** (built + ran + token-exact on real sm_110); (2a) **`marlin-nvfp4` = KERNEL-RUNTIME-VERIFIED (numerics-gated) on sm_110 as of 2026-08-11 — explicitly NOT model-level, see the Marlin evidence block at left; it does NOT upgrade `CLAIM-CUDA-SM110-RUNTIME`, whose token-exact bar remains met only by the portable bf16 path**; (2b) fp8/CUTLASS/FA2 + the fp4-mma/cutlass-nvfp4 cells = DERIVED/NOT-YET (still EMPTY for sm_110; a cutlass-backed kernel campaign, and cutlass is not even installed on Thor); (3) same-family siblings inherit the portable bodies but have NOT independently run. Precise scope: RUNTIME-VERIFIED applies ONLY to what actually RAN (portable bf16 Llama-1B greedy), NOT the whole arch | `CLAIM-CUDA-SM110-RUNTIME` | -| `BACKEND-CUDA-SM120` | CUDA compute 12.0 (consumer Blackwell, RTX 5070/5080/5090 class) | `>=12.8` | `CMakeLists.txt:105-118`; SM12 FP4 at `:940-970`; vLLM's own `FP4_SM120_ARCHS "12.0a;12.1a"` at `:951` and sm120 `SCALED_MM_ARCHS` at `:777` — i.e. upstream treats 12.0 and 12.1 as ONE capability set, which is why our FEATURE TABLE cells already named `12.0a` | Declared as DATA, no kernel code: all five `12.0a` cells of `VT_CUDA_FEATURE_TABLE` [CudaArchFeatures.cmake:216](../cmake/CudaArchFeatures.cmake#L216) resolved at [CMakeLists.txt:81](../CMakeLists.txt#L81); a documented target value on [`VLLM_CPP_CUDA_ARCHITECTURES`](../CMakeLists.txt#L37); runtime dispatch keyed on the Blackwell FAMILY not on GB10 — the fp4 tactic predicate is `caps.sm_major == 12` [cuda_matmul_nvfp4.cu:2653](../src/vt/cuda/cuda_matmul_nvfp4.cu#L2653), attention priority is `LookupAttnPriority(major)` [cuda.cpp:69](../src/vllm/platforms/cuda.cpp#L69); discrete/UMA classification [cuda_backend.cu:283](../src/vt/cuda/cuda_backend.cu#L283) requires both pageable access and `cudaDevAttrIntegrated`, which is exactly the discrete-sm_120 case. **Zero kernel/model/runner edits were required** | **BUILD evidence (dgx, nvcc 13.0, `~/work/sm120a`, 2026-07-22):** single-arch `120a` and the same-family fat `120a;121a` both configure with fp4-mma / cutlass-nvfp4 / cutlass-fp8 / marlin-nvfp4 / fa2 all `ENABLED for [120a]` (no missing-tactic warning); the fat build compiles and links `-Werror` **0 warnings** and `cuobjdump -lelf libvllm.a` shows **38 TUs carrying real `sm_120a` cubins** alongside 38 `sm_121a`, the fp4-MMA TU among them. Resolution is asserted, not eyeballed: [CudaArchFeaturesTest.cmake](../cmake/CudaArchFeaturesTest.cmake) is 35 hard expectations (`cmake -P`, no GPU) wired into CI, mutant-checked. **RUNTIME evidence exists only for the loader/backend leaf, from an external RTX 5070 Ti:** CUDA attribute gate `tests/vt/test_cuda_backend.cpp:153-172` asserts `UnifiedMemory == pageable && integrated` and passes 3/3; plain Qwen3.5-4B direct-load gate `tests/vllm/models/test_qwen35_plain_weights.cpp:162-196` passes 1664/1664; root `/tmp/qwen35-transplant-4b-647a2a6a` proves direct ON peak/stable PSS **2.606/0.727 GiB** vs OFF **8.559/8.558 GiB** with 128/128 paired output IDs. **NO gate model, kernel suite or benchmark has ever executed on sm_120 hardware here** | [CUDA inventory](specs/cuda-architecture-inventory.md); [arch additivity §W8](specs/cuda-arch-additivity.md) | `ACTIVE` - **claimed and BUILD-SUPPORTED; this is NOT runtime support** — the same sense in which the Metal row is claimed for a gated skeleton. Configure/compile/feature-resolution/SASS-emission are proven and the additive contract is proven (zero kernel edits); the loader/backend leaf has external 5070 Ti evidence. Everything else is UNTESTED because no sm_120 board exists here — a green fatbinary link is not execution evidence. Also: Triton AOT is single-arch by cubin nature, so a fat build needs `VLLM_CPP_TRITON=OFF` and an sm_120-only Triton build needs `triton_aot_vendored/sm_120a/` regenerated ON the card (deliberately absent; no cubins faked) | `CLAIM-CUDA-SM120-BRINGUP` | +| `BACKEND-CUDA-SM120` | CUDA compute 12.0 (consumer Blackwell, RTX 5070/5080/5090 class) | `>=12.8` | `CMakeLists.txt:105-118`; SM12 FP4 at `:940-970`; vLLM's own `FP4_SM120_ARCHS "12.0a;12.1a"` at `:951` and sm120 `SCALED_MM_ARCHS` at `:777` — i.e. upstream treats 12.0 and 12.1 as ONE capability set, which is why our FEATURE TABLE cells already named `12.0a` | Declared as DATA, no kernel code: all five `12.0a` cells of `VT_CUDA_FEATURE_TABLE` [CudaArchFeatures.cmake:216](../cmake/CudaArchFeatures.cmake#L264) resolved at [CMakeLists.txt:81](../CMakeLists.txt#L81); a documented target value on [`VLLM_CPP_CUDA_ARCHITECTURES`](../CMakeLists.txt#L37); runtime dispatch keyed on the Blackwell FAMILY not on GB10 — the fp4 tactic predicate is `caps.sm_major == 12` [cuda_matmul_nvfp4.cu:2653](../src/vt/cuda/cuda_matmul_nvfp4.cu#L2653), attention priority is `LookupAttnPriority(major)` [cuda.cpp:69](../src/vllm/platforms/cuda.cpp#L69); discrete/UMA classification [cuda_backend.cu:283](../src/vt/cuda/cuda_backend.cu#L283) requires both pageable access and `cudaDevAttrIntegrated`, which is exactly the discrete-sm_120 case. **Zero kernel/model/runner edits were required** | **BUILD evidence (dgx, nvcc 13.0, `~/work/sm120a`, 2026-07-22):** single-arch `120a` and the same-family fat `120a;121a` both configure with fp4-mma / cutlass-nvfp4 / cutlass-fp8 / marlin-nvfp4 / fa2 all `ENABLED for [120a]` (no missing-tactic warning); the fat build compiles and links `-Werror` **0 warnings** and `cuobjdump -lelf libvllm.a` shows **38 TUs carrying real `sm_120a` cubins** alongside 38 `sm_121a`, the fp4-MMA TU among them. Resolution is asserted, not eyeballed: [CudaArchFeaturesTest.cmake](../cmake/CudaArchFeaturesTest.cmake) is 35 hard expectations (`cmake -P`, no GPU) wired into CI, mutant-checked. **RUNTIME evidence exists only for the loader/backend leaf, from an external RTX 5070 Ti:** CUDA attribute gate `tests/vt/test_cuda_backend.cpp:153-172` asserts `UnifiedMemory == pageable && integrated` and passes 3/3; plain Qwen3.5-4B direct-load gate `tests/vllm/models/test_qwen35_plain_weights.cpp:162-196` passes 1664/1664; root `/tmp/qwen35-transplant-4b-647a2a6a` proves direct ON peak/stable PSS **2.606/0.727 GiB** vs OFF **8.559/8.558 GiB** with 128/128 paired output IDs. **NO gate model, kernel suite or benchmark has ever executed on sm_120 hardware here** | [CUDA inventory](specs/cuda-architecture-inventory.md); [arch additivity §W8](specs/cuda-arch-additivity.md) | `ACTIVE` - **claimed and BUILD-SUPPORTED; this is NOT runtime support** — the same sense in which the Metal row is claimed for a gated skeleton. Configure/compile/feature-resolution/SASS-emission are proven and the additive contract is proven (zero kernel edits); the loader/backend leaf has external 5070 Ti evidence. Everything else is UNTESTED because no sm_120 board exists here — a green fatbinary link is not execution evidence. Also: Triton AOT is single-arch by cubin nature, so a fat build needs `VLLM_CPP_TRITON=OFF` and an sm_120-only Triton build needs `triton_aot_vendored/sm_120a/` regenerated ON the card (deliberately absent; no cubins faked) | `CLAIM-CUDA-SM120-BRINGUP` | | `BACKEND-CUDA-SM121` | CUDA compute 12.1 / current `121a` gate | `>=12.8,<13` globally; family via `f` on 13 | `CMakeLists.txt:105-118`; SM12 FP4 at `:940-970` | default target [CMakeLists.txt:37](../CMakeLists.txt#L37); runtime [cuda_backend.cu:20](../src/vt/cuda/cuda_backend.cu#L20); SM12 fast paths now resolved per arch through the FEATURE TABLE [CMakeLists.txt:68](../CMakeLists.txt#L68), [CudaArchFeatures.cmake:1](../cmake/CudaArchFeatures.cmake#L1); the sole registered runtime tactic [cuda_arch_tactics.h:1](../src/vt/cuda/cuda_arch_tactics.h#L1) | [CUDA backend tests](../tests/vt/test_cuda_backend.cpp#L31); [27B gate](../tests/parity/test_qwen27_paged_engine.cpp#L110); [35B gate](../tests/parity/test_qwen36_paged_engine.cpp#L78); [ledger trace](parity-ledger.md#L284) | [CUDA inventory](specs/cuda-architecture-inventory.md); [arch additivity](specs/cuda-arch-additivity.md) | `PARTIAL` - gate workload only, not full family coverage | - | ## CUDA component target rules @@ -288,7 +288,7 @@ the rest are `SPIKE`. | ID | Item | Upstream | Our code | Tests/evidence | Spike/spec | State | Owner | |---|---|---|---|---|---|---|---| -| `BACKEND-DISTRIBUTED-COMM` | The unifying `vt::Communicator` / process-group abstraction — rank/world_size + AllReduce(sum/max/min/prod)/AllGather/Send/Recv, stream-ordered (each takes a `Queue&`). **W1 LANDED**: abstraction (`include/vt/communicator.h`) + a CPU in-process multi-rank transport (`src/vt/communicator.cpp`, N ranks = N host threads over one barrier+staging+mailbox) proven by `tests/vt/test_communicator.cpp` (2/4-rank AllReduce-sum + AllGather exact on every rank, Send/Recv rendezvous, RED-verified; 8 cases/50 assertions). `world_size==1` ⇒ every collective a byte-identical no-op (asserted). **W2 LANDED**: collectives now ROUTE through `OpProvider`/`OpId` (`kAllReduce`/`kAllGather`/`kSend`/`kRecv`, keyed on the queue's DeviceType) — the CPU in-process reduce registered on kCPU (`test_communicator` still 50/50 through the OpId path), the NCCL provider on kCUDA. W2+ residuals: RDMA/TCP (Spark), MLX-ring (kMETAL) transports | vLLM `device_communicators/base_device_communicator.py:147` (DeviceCommunicatorBase interface, the port template) + `distributed/parallel_state.py:358` (GroupCoordinator dispatch; world_size==1 bypass :638) | LANDED: `include/vt/communicator.h` + `src/vt/communicator.cpp` (sibling of `vt::Queue` `include/vt/device.h:50`); OpId routing via `include/vt/op_provider.h:108` (`OpId::kAllReduce/…`); stream-order hooks reused `include/vt/backend.h:87-104` | CPU exact-gate (`test_communicator`, 50/50, via OpId path) | [scale-out spike](specs/scale-out-distributed.md) | `ACTIVE` | `CLAIM-SCALE-OUT-W2` | +| `BACKEND-DISTRIBUTED-COMM` | The unifying `vt::Communicator` / process-group abstraction — rank/world_size + AllReduce(sum/max/min/prod)/AllGather/Send/Recv, stream-ordered (each takes a `Queue&`). **W1 LANDED**: abstraction (`include/vt/communicator.h`) + a CPU in-process multi-rank transport (`src/vt/communicator.cpp`, N ranks = N host threads over one barrier+staging+mailbox) proven by `tests/vt/test_communicator.cpp` (2/4-rank AllReduce-sum + AllGather exact on every rank, Send/Recv rendezvous, RED-verified; 8 cases/50 assertions). `world_size==1` ⇒ every collective a byte-identical no-op (asserted). **W2 LANDED**: collectives now ROUTE through `OpProvider`/`OpId` (`kAllReduce`/`kAllGather`/`kSend`/`kRecv`, keyed on the queue's DeviceType) — the CPU in-process reduce registered on kCPU (`test_communicator` still 50/50 through the OpId path), the NCCL provider on kCUDA. W2+ residuals: RDMA/TCP (Spark), MLX-ring (kMETAL) transports | vLLM `device_communicators/base_device_communicator.py:147` (DeviceCommunicatorBase interface, the port template) + `distributed/parallel_state.py:358` (GroupCoordinator dispatch; world_size==1 bypass :638) | LANDED: `include/vt/communicator.h` + `src/vt/communicator.cpp` (sibling of `vt::Queue` `include/vt/device.h:131`); OpId routing via `include/vt/op_provider.h:108` (`OpId::kAllReduce/…`); stream-order hooks reused `include/vt/backend.h:87-104` | CPU exact-gate (`test_communicator`, 50/50, via OpId path) | [scale-out spike](specs/scale-out-distributed.md) | `ACTIVE` | `CLAIM-SCALE-OUT-W2` | | `BACKEND-DISTRIBUTED-TP` | Tensor parallel (intra-node multi-GPU) — sharded Column/Row/QKV linears, vocab-parallel embed + LM head, attention-head split, MoE expert-parallel; all-reduce after o_proj/MLP-down and the EP combine. **W2 LANDED (CPU-gated)**: the `TensorParallel`/`TpShard`/`TpAllReduceSum` wiring (`include/vllm/model_executor/models/tensor_parallel.h`) threaded into the Qwen3-dense forward (o_proj all-reduce `dense_attn_block.h`, MLP-down `qwen3.cpp`) + the MergedColumn shard at the loader chokepoint (`dense_weight_loaders.h`), proven by `tests/vt/test_tp_forward.cpp` — sharded-matmul + RowParallel all-reduce **== the unsharded tp=1 forward** over the W1 CPU communicator (RED-verified: dropping the all-reduce fails 24 assertions). `tp_size==1`/nullptr ⇒ every helper a byte-identical no-op (asserted). RESIDUAL (HW-gated, no ≥2-GPU box): QKV head-aware KV replication, vocab/LM-head + MoE-EP sharding, and a real TP-2 GPU forward | vLLM `layers/linear.py:418` (Column, out-dim shard) / `:1612` (Row, all-reduce :1766) / `:1021` (QKV heads :1074) + `vocab_parallel_embedding.py:198` + `fused_moe/expert_map_manager.py:22` | LANDED `include/vllm/model_executor/models/tensor_parallel.h`; seams `dense_attn_block.h` (o_proj all-reduce) + `qwen3.cpp` (MLP-down); weight chokepoint `dense_weight_loaders.h:131` (column shard) | CPU multi-rank TP gate (`test_tp_forward`, 60/60, RED-verified) | [TP end-to-end spike](specs/tensor-parallelism-spike.md) (2026-08-08, task #287 — at-pin S1 inventory + landed-vs-claimed audit + TP-W0..W7 plan; TP-W1..W4+W7 CPU-completable) + [scale-out spike](specs/scale-out-distributed.md) | `READY` | - | | `BACKEND-DISTRIBUTED-PP` | Pipeline parallel — PP stage split (`PPMissingLayer` analogue) + inter-stage `IntermediateTensors` send/recv over the comm layer + multi-worker executor fan-out | vLLM `models/utils.py:785` (PPMissingLayer) / `:798` (make_layers) + `distributed/utils.py:127` (get_pp_indices) + `parallel_state.py:957` (send_tensor_dict) | fan-out seam `src/vllm/v1/executor/executor.cpp:7-34` (direct single-worker call today) | - | [scale-out spike](specs/scale-out-distributed.md) | `SPIKE` | `CLAIM-SCALE-OUT-SPIKE` | | `BACKEND-DISTRIBUTED-DP` | Data parallel — N independent engine replicas over the SAME weights + a DP coordinator (global "request wave" so all DP ranks step together) + a per-step token-count all-reduce; DP×EP is the large-scale DeepSeek serving topology (DP-replicated attention + EP-sharded experts). NOT part of `world_size` (DP is outside: `world_size_across_dp = world_size × DP`) | vLLM `v1/engine/coordinator.py:23` (`DPCoordinator`, wave :33-56) + `v1/worker/dp_utils.py:164` (`coordinate_batch_across_dp`; per-step `num_tokens_across_dp` all-reduce :53) + group `distributed/parallel_state.py:1866` + flags `config/parallel.py:129-145` | reuses W1 `Communicator::AllReduce` for the token-count sync; NEW engine-replica executor + coordinator; **depends on the multi-worker executor (`executor.cpp:7-34`)** | - | [parallelism-modes spike](specs/parallelism-modes.md) | `SPIKE` | `CLAIM-PARALLELISM-MODES-SPIKE` | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index b26a22dc33..93880222d8 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -28723,13 +28723,12 @@ killed every measured leg mid-load on the previous attempt. 0.774 GiB on disk in bf16 and 1.547 GiB resident, because `qwen3_vl.cpp` widens it to host f32 — [#1359](https://github.com/mudler/vllm.cpp/issues/1359), which also affects - the Qwen3.6-27B path. #1359's Qwen3-VL half has since LANDED, and the - 2026-08-28 rerun recorded later in this file MEASURED the consequence: - 826916864 B = 0.770 GiB, **0.499x this figure**. The HALVING IS CORRECT - rather than a regression — the flag now frees the tower the checkpoint ships - instead of the tower plus our widening. The figure recorded here stands - unaltered as what the run at `41ab550b9` measured, and is superseded for - current behaviour; `muse-glimmer-30b` still widens, blocked on + the Qwen3.6-27B path. #1359's Qwen3-VL half has since LANDED, so this leg + rerun should read about 0.774 GiB rather than 1.542, and that HALVING IS + CORRECT rather than a regression — the flag now frees the tower the + checkpoint ships instead of the tower plus our widening. The figure recorded + here stands as what the run at `41ab550b9` measured; `muse-glimmer-30b` still + widens, blocked on [#2166](https://github.com/mudler/vllm.cpp/issues/2166). 2. **Load-time residency, not a served request.** `ForwardQwen3VLForConditionalGeneration` `VT_CHECK`s `input.mm.has_value()`, diff --git a/.agents/completed/issue-index.md b/.agents/completed/issue-index.md index 199a10828f..e86d948171 100644 --- a/.agents/completed/issue-index.md +++ b/.agents/completed/issue-index.md @@ -554,286 +554,114 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1563](https://github.com/mudler/vllm.cpp/issues/1563) | `GATE-SQUASH-SEPARATOR` | **A markdown `---` horizontal rule anywhere in a pull request body silently voids the trailer block, and `check-commit-trailers.py` blames the trailers instead of the framing.** Found 2026-08-21 writing the body for PR [#1550](https://github.com/mudler/vllm.cpp/pull/1550) ([#1542](https://github.com/mudler/vllm.cpp/issues/1542)). `parsed_trailers()` shells out to git's trailer parser, and **git treats a line of exactly `---` as the start of the patch section**, so everything after the first one is not part of the message and a trailer block below it is invisible. Reproduced with no repository state: a body of `subject / prose / --- / more prose / FOLLOWING_AGENTS_PROTOCOL / the three trailers` reports `[trailers] Following-Agents-Protocol must appear exactly once` and `[attribution] AI-Assisted must appear exactly once`; `sed -i '/^---$/d'` on that same file reports `OK: commit trailer contract`, and the `---` is the only difference. **The MESSAGE is the defect, not only the behaviour**: `Following-Agents-Protocol` appears EXACTLY ONCE in the body while the checker says it must appear exactly once, so a reader counts occurrences, finds one, counts again and dumps bytes before thinking to test the parser's own framing. `_strict_errors` already computes `_paragraphs(body)[-1]` correctly as the three trailers verbatim, so the checker holds the information needed to say "the trailer paragraph is present but git could not parse it; a `---` line at line N ends the message". Worse, the neighbouring `FOLLOWING_AGENTS_PROTOCOL must appear exactly once as a separate paragraph before the trailer paragraph` check stays SILENT, so the two errors that fire both point away from the cause. **Beyond one confusing message**: the repository sets `squash_merge_commit_message = PR_BODY`, so the body IS the landed commit message, and a body carrying a `---` lands a commit whose trailers `git interpret-trailers` cannot see, on a branch that is never force-pushed. Same permanent-damage shape AGENTS.md records for the `---------` separator GitHub wrote under `COMMIT_MESSAGES`, arriving from the AUTHOR side rather than the forge side. `scripts/agent-pr-body.py --pr ` DOES catch it and caught it here before the merge; the exposure is a body never passed through that command, which AGENTS.md notes is not a gate and cannot be one because it reaches the network, while the CI guard reads the frozen `pull_request` payload and so does not re-read a body edited after the final push. NOT FIXED HERE: it changes a checker's semantics and its message, so under `## Changing the rules or a checker` it needs its own row, a red-before test and green-after evidence. Two candidate repairs, neither chosen: name the `---` line, or strip patch-section framing before parsing so a markdown rule is inert -- the second changes what the contract accepts and is the larger decision. Suggested minimum: `tests/scripts/test_check_commit_trailers.py` gains a case pinning the reproduction above | bug | | [#1575](https://github.com/mudler/vllm.cpp/issues/1575) | `SPEC-DFLASH2` | **`build-newest-gcc` has been RED on `main` since `5702d8f83`, and it is a recurrence of the class `tests/support/process_id.h` was created to close, not a missing include.** `tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547` calls `::getpid()` while including nothing that declares it; the lane builds in a `gcc:16` container where libstdc++ no longer supplies the transitive ``, so it fails to COMPILE (`error: '::getpid' has not been declared; did you mean 'getpt'?`) while every local toolchain at gcc 13 stays green. The seam's own header predicts this: it says the spelling "was fixed once in three files and came back in five more, because each new loader test copies the temp-directory helper from the last one", which is exactly what W5's GGUF test did. The convention is otherwise unanimous — every other file in the tree naming `getpid` either includes `` or uses the helper, and the sibling `tests/vllm/entrypoints/test_dspark_draft_routing.cpp` uses the helper AND carries a comment warning not to reintroduce this class. ATTRIBUTED, not inferred: `build-newest-gcc` is absent from the baseline's failed list at `92406c620` and present at `5702d8f83`, the commit that added the file. MEASURED red-first in the lane's own container at base `947e5f648`, file sha256 `83bba319…`: unmodified `RED_RC=1` with the error byte-identical to CI; with the seam applied `GREEN_RC=0`, `compile_err=0`, `git diff --stat` confirming the edit applied; local gcc 13 `GCC13_RC=0`, so no regression on the shipped toolchain. FIXED IN FLOW with the portable spelling (`#include "support/process_id.h"`, `vllm_test::ProcessId()`) rather than ``, which would work on POSIX but re-copies the idiom the helper centralises and does not compile on MSVC. Found while measuring [#1464](https://github.com/mudler/vllm.cpp/issues/1464) at `origin/main` | bug | | [#1454](https://github.com/mudler/vllm.cpp/issues/1454) | `SPEC-MTP-GGUF` | **`test_qwen3_5_gguf_mtp.cpp` reported `Status: SUCCESS!` with `assertions: 0` on every CI run, and its one arithmetic guarantee was a tautology.** Both cases opened `if (path == nullptr) return;` on `VLLM_MTP_GGUF_MODEL`, and a bare `return` from a doctest case is a PASS: re-derived on a clean Release build at `947e5f648`, unset, the file printed `test cases: 2 \| 2 passed \| 0 failed \| 0 skipped`, `assertions: 0`, `Status: SUCCESS!`, exit 0, and printed nothing else. The variable is set nowhere in `.github/workflows/`, so that was the state of every run. Second defect in the same file: the comment at `:52` stated `num_hidden_layers + depth == block_count` and the line under it asserted `CHECK(c.num_hidden_layers > 0)`, true of every valid model. MEASURED, not argued: mutating `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:889` to `c.num_hidden_layers = block_count;` compiled clean and left the file at 2/2 cases, 0 assertions, `SUCCESS!`, exit 0. FIXED IN FLOW. The invariant is now pinned **HERMETICALLY** on KV-only synthetic GGUFs carrying no weight bytes, so CI checks it every run rather than never - 65/1 (the shipped Qwen3.8-27B pair), 25/1 (the Qwen3.5-2B reference this suite was developed against) and 28/3, the third arm separating `- nextn` from `- 1` - plus a head-less arm asserting the key is NOT published, which is the half `NumMtpLayers` cannot express because it answers 1 for an absent key. The two env-gated cases stay, now skipping with a `MESSAGE` naming the variable as `test_gguf_mmproj_reach.cpp` does, and the live one re-derives the invariant from the file's own `block_count` kv. Unset 4 cases / 18 assertions / `SUCCESS!` / rc 0; live on `Qwen3.8-27B-Q4_K_M.gguf` 4 / 38 / `SUCCESS!` / rc 0. Both mutants now red (9/18 and 5/18, exit 1), compiled clean, restored against a pre-taken sha256. **The production line is CORRECT and was not touched**: `block_count - nextn` landed `1a4db5c3c`, the `mtp_num_hidden_layers` republication `493327b4e`. Related but distinct: [#821](https://github.com/mudler/vllm.cpp/issues/821) W2 (`0adeb8b0e`) pins the same arithmetic for the 27B artifact on a committed manifest in `tests/vllm/models/test_qwen38_27b_gguf_manifest.cpp`, and that gate DOES catch both mutants - so the invariant was not globally unpinned, it was unpinned in this row's own file | bug | -| [#1434](https://github.com/mudler/vllm.cpp/issues/1434) | `GATE-DOC-CHECKPOINT-STATES` | **`scripts/check-doc-checkpoint.py` could not see `PARTIAL`, so 118 state cells could move with no gate observing them.** `STATES` (`:56-66`) is the whole definition of what a lifecycle state IS for the gate that enforces AGENTS.md's `docs/STATUS.md` / `docs/BENCHMARKS.md` / spec `## Now` triple, and `row_states` drops any row it cannot match. `lifecycle_moves` and `moved_rows` then iterate the AFTER map, so leaving the matched set is silent by construction. Re-derived at `947e5f648` (the report measured `63d87805c`): `PARTIAL` **118** cells and `ANCHOR-BACKFILL` **73**, against `DONE` 77 and `BLOCKED` 9 — `PARTIAL` is the second most used state in the matrices and the gate was blind to it. Over the seven tables `ROW_TABLES` actually reads, the resolved population goes from **153 rows to 226**, a 47.7 % widening. Two of the transitions the report names behave differently from its description, measured with scratch commits at `947e5f648` on an unmodified checker: `READY -> PARTIAL` rc **0** and `PARTIAL -> READY` rc **0** are the real blind spots, while the report's suggested `PARTIAL -> ACTIVE` already reds — by accident, reporting **`added as ACTIVE`** for a row that has existed for months, because it is absent from the BEFORE map. FIXED IN FLOW for `PARTIAL` only. **`ANCHOR-BACKFILL` is deliberately excluded**: `.agents/feature-matrix.md:14-17` defines it as a property of the RECORD (*a legacy implemented row without exact code, test and real-spec anchors*), `docs/STATUS.md` carries no such term and would have nothing true to write on a `DONE <-> ANCHOR-BACKFILL` move, and `REQUIRED["lifecycle"]` cannot demand the spec's `## Now` alone — so admitting it would demand a public-document edit with nothing to say, which is the exact shape `check-doc-checkpoint.py:4-17` records as the reason the file was rewritten (16 of 20 red CI runs, six hardcoded escape hatches). One row's resolved state moves and the move is a REPAIR: `KV-BLOCK-POOL` says `` `PARTIAL` (not `DONE`) `` in its prose and the last-match heuristic believed the parenthesis, resolving `DONE`. No pinned counter moves — `check-gate-commands.py` has its own `GATED_STATES` and `RUNNABLE_BASELINE` is keyed on matrix rows, `UNOWNED_HIGH_WATER` is unmoved because this row names an owner, and no matrix row or public document changes — which was measured, not assumed, because this is the [#1376](https://github.com/mudler/vllm.cpp/issues/1376) ratchet shape. Remainder listed under `## Owed` in [doc-checkpoint-lifecycle-states.md](../specs/doc-checkpoint-lifecycle-states.md): `ANCHOR-BACKFILL` moves, `.agents/sglang-matrix.md` never entering `ROW_TABLES`, a row that leaves the matched set entirely, and a new row added directly as `PARTIAL` | bug | -| [#1549](https://github.com/mudler/vllm.cpp/issues/1549) | `LTX25-DIT-ATTN-FLASH` | **One LTX-2.5 DiT forward costs 47.84 s on GB10 because the DiT self-attention never opted into a fast attention op.** `src/vllm/model_executor/models/ltx2_device.cpp:421` calls `vt::Attention`, which on CUDA is `AttentionKernel` (`src/vt/cuda/cuda_ops.cu:1463`) -- the kernel whose own header at `:1456-1459` calls itself "Correctness-grade (M0.9)": one 256-thread block per (query, head), a 256-wide shared-memory tree reduction per key, no K/V tiling. At `768x448/49f` (2352 tokens) that is 75,264 blocks x 2352 keys = 1.77e8 block-key iterations per call x 48 layers. MEASURED 47.84 s mean / 47.91 s median, n=119, spread 5.8%, from the engine's own `last=` lines (`render_phase_log.cpp:388` via `ltx2_video.cpp:4048-4054`), never from the governor, which has reported 1.00 s, 69.1 s, 162 s and 396.9 s for this one quantity. Attribution is arithmetic, not assertion: `.agents/specs/multimodal-speed.md:24-26` measures this same kernel on this same box at 5.70 ns per block-key iteration, and 1.77e8 x 5.70 ns x 48 = **48.4 s against the measured 47.84 s, a 1% match**. WHY IT WAS MISSED, which is the reusable part: `kAttention` is deliberately frozen on the naive kernel so text decode stays byte-identical (`cuda_ops.cu:3120-3122`), and the fast kernels are SEPARATE OPS each caller must opt into BY NAME. No automatic selection, no fallback notice. A model that never opts in gets correct output at ~500x the cost with no warning anywhere -- goldens pass, no refusal fires, and `GetOpProviderStats` counts the naive selection as the success it is. FIXED IN FLOW by routing the self-attention to `vt::AttentionDenseFlash`. Two things found doing it and filed rather than folded in: [#1551](https://github.com/mudler/vllm.cpp/issues/1551) (FA-2 refuses head_dim 128) and [#1552](https://github.com/mudler/vllm.cpp/issues/1552) (the same defect shape at every other `vt::Attention` caller). Spec [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) | bug | -| [#1551](https://github.com/mudler/vllm.cpp/issues/1551) | `LTX25-DIT-ATTN-FLASH` | **`vt::AttentionDenseFa2` refuses anything but head_dim 64, so LTX-2.5's head_dim-128 DiT cannot reach tensor cores.** The guard is `src/vt/cuda/cuda_flash_attn_fa2.cu:557-560` plus the dispatch test `query.shape[2] == 64` at `src/vt/cuda/cuda_ops.cu:3396-3399`; everything else falls through to `AttentionDenseFlash`, which is a scalar warp-per-query online-softmax recurrence with shared-memory K/V tiling -- correct, tiled, and still not `mma.sync`. LTX's video stream is 32 heads x head_dim 128 (`include/vllm/model_executor/models/ltx2.h:124-125`). Reaching the vendored FA-2 path needs an extra `run_mha_fwd_` instantiation and the guard widened; `cuda_ops.cu:3375-3377` records the current narrowness as deliberate, so this is deferred cost and not oversight. Prize is BOUNDED, not measured: `.agents/specs/multimodal-speed.md` §16 puts warp-to-flash at 1.04x at 784 tokens and §14 at 1.82x at 1500, both the same scalar recurrence; what tensor cores buy over it at 2352 tokens is unmeasured, and this issue owes the measurement before it owes the port. Numerics caveat that must not be lost: FA-2 is NOT bit-identical (`include/vt/ops.h:2995-2997`, `mma.sync` reassociates both QK^T and P.V), and a diffusion model has no token gate, so a pixel-level comparison has to be designed first. NOT fixed in flow -- explicitly out of scope for #1549. Owner: row `LTX25-DIT-ATTN-FLASH`, under `## Owed` in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) | feature | -| [#1552](https://github.com/mudler/vllm.cpp/issues/1552) | `LTX25-DIT-ATTN-FLASH` | **A caller that never opts into a fast attention op is never told, so sweep every remaining `vt::Attention` call site.** The generalisation of [#1549](https://github.com/mudler/vllm.cpp/issues/1549), filed separately because the mechanism is not LTX-specific. `kAttention` is frozen on the naive kernel so text decode stays byte-identical (`src/vt/cuda/cuda_ops.cu:3120-3122`) -- correct, and untouched here. The consequence is that `kAttentionDenseFast` / `kAttentionDenseFlash` / `kAttentionDenseFa2` are separate ops each caller must name, with no shape routing and no fallback notice, so a caller that never opts in gets correct output at up to ~500x the cost and NOTHING detects it: the goldens pass because the output is right, no refusal fires because the op is registered, and `GetOpProviderStats` counts the naive selection as a success because it is one. The only symptom is a wall clock the model may have no gate for. OWES: (1) every non-decode `vt::Attention` call site enumerated with its head_dim, sequence length and dense/non-causal eligibility; (2) per eligible site, either a routing change with its own reachability proof and numerics gate or a recorded reason to stay; (3) a decision between leaving it caller-opt-in, warning once on a large-token `kAttention` selection on CUDA, or shape-routing `kAttention` itself -- only the third removes the failure mode and only the third risks the byte-identity guarantee, so it needs its own spec. Same class as AGENTS.md "Nothing lands dead", inverted: there a capability lands unreached, here a FASTER capability lands unreached and the slow one is correct enough that nobody looks. NOT fixed in flow. Owner: row `LTX25-DIT-ATTN-FLASH`, under `## Owed` in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) | bug | -| [#1612](https://github.com/mudler/vllm.cpp/issues/1612) | `LTX25-DIT-ATTN-FLASH` | **The LTX-2.5 DiT attention swap has NO numeric or pixel comparison at production geometry.** #1549 moved the device DiT self-attention from `vt::Attention` to `vt::AttentionDenseFlash`, and the two are explicitly NOT bit-identical on CUDA: the warp kernel groups the head_dim partial sums across 32 lanes where the naive kernel uses a 256-thread block, so the same f32 online softmax associates differently (`.agents/specs/ltx25-dit-attn-flash.md` §5). The ONLY numeric gate that exists is `test_ltx2_device`'s CUDA host-vs-device parity case -- video `8.9407e-08`, audio `4.47035e-08` against a committed `2e-5` on `dgx:gpu0` -- and it runs the REDUCED-DIMENSION fixture, so it bounds the ARITHMETIC change (a length-D sum reassociated) and says nothing about the change at head_dim 128 with 2352 keys over 48 layers, which is what production runs. A diffusion render has no token gate to fall back on, which is the mechanism every other model here leans on; the substitute is a frame comparison against a baseline render, and the flash arm's run (`20260821T092516Z`) was interrupted at forward 20 when the `rc` worker was lost and **wrote no frames at all**, so nothing exists to compare even against the completed 49-frame 768x448 baseline already on the NAS. OWES: (1) a COMPLETING flash-arm render at `768x448/49f` seed `20260820` on the recorded checkpoints; (2) a pixel comparison criterion DESIGNED before the numbers are read, because the arms are not bit-identical by construction and "the frames differ" is not a finding -- the same design question [#1551](https://github.com/mudler/vllm.cpp/issues/1551) owes for the larger FA-2 divergence, so one answer should serve both; (3) failing those, a same-binary A/B render pair under one lease with `VLLM_LTX2_DIT_FLASH_ATTN` unset and `=0`, which yields the pixel pair and the speed pair at once -- `scripts/ltx25-dit-attn-flash-ab.sh` already runs both arms from one binary but caps each at 13 samples for the timing. NOT fixed in flow: it needs `dgx:gpu0` for two full renders and a criterion agreed first, which is a design question and not a repair. Owner: row `LTX25-DIT-ATTN-FLASH`, under `## Owed` in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) | bug | -| [#1584](https://github.com/mudler/vllm.cpp/issues/1584) | `KERNEL-ATTENTION-CROSS-CUDA` | **`OpProviderStats::declines` double-counts the FIRST decline of every process, and `include/vt/op_provider.h` states the opposite.** That header prescribes a pattern for a SHAPE-GATED provider -- cache the fallback pointer once, call `NoteOpDecline` per decline -- and says it keeps the count "exact" while removing the lookup cost, citing a decode run that declines ~21,500 times. But `GetOpFallback` does its own `slot.declines.fetch_add(1)` at `src/vt/op_provider.cpp:709`, and the prescribed pattern resolves the fallback through a function-local static, so the first decline takes TWO increments and every later one takes one. Live at `src/vt/metal/metal_mlx_provider.mm` `MlxFallback`:215-224 with `MlxMatmulKernel`:228 and `MlxMatmulBTKernel`:238 -- the site the header was written from -- and inherited by `src/vt/cuda/cuda_attention_cross.cu` `AttentionCrossBlockedCuda`, added by [#1555](https://github.com/mudler/vllm.cpp/issues/1555). **The count is small and the failure mode is not.** `declines` is the ONLY discriminator that separates a served call from a forwarded one for a provider that is always the SELECTED one, so it is what any routing gate must assert; a gate asserting `declines == N` therefore passes or fails on whether some earlier case in the same binary already forced the static to resolve. Measured in `tests/vt/test_ops_attention_cross.cpp`: green on a full run, and `doctest -tc=` on a single case reads 2. A test passing on the order of its file is not measuring what it claims. **NOT fixed in flow, and the reason is scope rather than difficulty.** Dropping the `fetch_add` at `:709` would leave `GetOpFallback` a pure lookup and `NoteOpDecline` the sole counter, but the callers that use `GetOpFallback` WITHOUT the caching pattern -- `src/vt/vulkan/vulkan_ops.cpp:950,1067,1488,1509` and `src/vt/tenstorrent/tenstorrent_ops.cpp:1341` -- would then stop counting entirely, so it changes what four backends report and needs its own row, a red-before test and green-after evidence per AGENTS.md `## Changing the rules or a checker`. A second candidate is a `GetOpFallbackUncounted()` so the two spellings differ by whether they count rather than by accident. **Worked around** by #1555, which warms the fallback static outside every counted window so its routing assertions are order-independent and exact. That makes those tests honest; it does not fix the seam. Listed under `## Owed` in [attention-cross-blocked.md](../specs/attention-cross-blocked.md) | bug | -| [#1598](https://github.com/mudler/vllm.cpp/issues/1598) | `SPEC-DRAFTER-CHAIN` | **A chain entry of `--speculative-config` called `draft_sample_method` and `rejection_sample_method` a typo, when this engine HONOURS both at the top level of the same document.** Measured at `31cefe631` + `e2a9e035d`: `{"vllm_cpp":{"drafter_chain":[{"method":"mtp","draft_sample_method":"greedy"}]}}` returned `unknown key "vllm_cpp.drafter_chain[0].draft_sample_method"` — the SAME message, modulo the name, that a real misspelling `num_speculatve_tokens` returns. Both keys are genuine `SpeculativeConfig` fields at the parity pin `555967922` (`Literal` aliases at `vllm/config/speculative.py:77,78`, field declarations at `:283` and `:216`), and `CheckValueGatedKey` value-gates both at the top level to upstream's own default, beside a chain, because the verify is engine-wide. Inside an ENTRY they were in neither `kHonouredKeys` nor `kUpstreamUnimplementedKeys`, so `CheckEntryKeys` fell them through to the typo branch — #1160's failure inverted, since #1160 split the classes precisely so that a key vLLM declares does not read as "unknown". Two prose claims were false with it: the function's own comment said it reused "the SAME three classes #1160 established" (it reused two), and `docs/SPECULATIVE-DECODING.md` said entry keys are refused "in the same two classes as above". **FIXED IN FLOW** in the `SPEC-DRAFTER-CHAIN` W1 repair. The correct class is neither honoured nor unimplemented — the keys ARE implemented, they are simply not per-drafter — so class 2 gets its own refusal naming the key, saying the engine honours it, and saying to spell it at the TOP LEVEL beside the chain. Class 2 is now one named `kEngineWideValueGatedKeys` set read by both admissions rather than two hand-written comparisons, which is how the split was lost. Red-first: the new subcase failed 4 assertions before the fix; mutating the class-2 branch out reddens it again. Found by a fresh review of `row/SPEC-DRAFTER-CHAIN-W1` | bug | -| [#1599](https://github.com/mudler/vllm.cpp/issues/1599) | `SPEC-DRAFTER-CHAIN` | **A non-string `model` is silently dropped, and the required-key message that follows names the key as MISSING when the user gave it.** Measured at `31cefe631` + `e2a9e035d`: `{"vllm_cpp":{"drafter_chain":[{"method":"mtp","model":123}]}}` parses, with `entry[0].draft_model_path` empty; `{..."method":"dflash","model":123}` refuses with `requires a "model" key naming the draft checkpoint`, pointing away from the actual mistake. `src/vllm/config/speculative.cpp:230-232` reads the key only when `is_string()` already holds, so every other JSON type takes the silent branch. **INHERITED, not introduced**: the top level has the identical shape at `:526-528` and behaves identically — `{"method":"dflash","model":123}` reports the key as missing too — so the chain entry is faithful to the landed contract and repairing only the entry would leave two spellings of one rule disagreeing. The repair belongs to both together: judge PRESENCE, then TYPE, and say "must be a string" when a value of the wrong type was given, which is #1160's polarity applied to value type instead of key presence. NOT fixed in `SPEC-DRAFTER-CHAIN` W1: it moves a landed top-level refusal that other suites assert on. Listed under `## Owed` in [drafter-chain.md](../specs/drafter-chain.md). Found by a fresh review of `row/SPEC-DRAFTER-CHAIN-W1` | bug | -| [#1600](https://github.com/mudler/vllm.cpp/issues/1600) | `SPEC-DRAFTER-CHAIN` | **A misspelled or mis-cased `vllm_cpp` on a chain-only document is answered with `a string "method" is required` — the one key such a document must not have.** Measured at `31cefe631` + `e2a9e035d`: `{"VLLM_CPP":{"drafter_chain":[...]}}`, a bare `{"drafter_chain":[...]}` and `{"vllm_cp":{...}}` all return that message. Each document names every speculator it wants; the engine demands the one key that D7 makes MUTUALLY EXCLUSIVE with a chain, and the engine says so itself when the spelling is right — so the user is told to add the key that would then be refused. Mechanism: `has_chain = doc.contains("vllm_cpp")` is false, so the method requirement at `src/vllm/config/speculative.cpp:353-379` fires before the key-admission loop at `:400-429`, which is where the unknown name would have been reported by name with the accepted list. **The ORDERING is inherited and deliberate** — #1160 put the method check first so an unsupported method is the first error a user sees — but the document CLASS that hits it is invented by this wave: before the chain existed, a document with no `method` was simply an incomplete vLLM document and the message was right. `.agents/specs/drafter-chain.md` D9 argues at length that the user must not be misled here, and this is the one shape where the landed code misleads. Two candidate repairs, neither chosen: report an unadmitted top-level key BEFORE the method requirement when the document carries no `method` at all, or judge the extension key case-insensitively for the DIAGNOSTIC only. Both move a landed error ordering that `tests/vllm/config/test_speculative_unknown_keys.cpp` and this row's own regression case assert on, so it needs its own row, a red-before test and green-after evidence. Listed under `## Owed` in [drafter-chain.md](../specs/drafter-chain.md). Found by a fresh review of `row/SPEC-DRAFTER-CHAIN-W1` | bug | -| [#1601](https://github.com/mudler/vllm.cpp/issues/1601) | `SPEC-DRAFTER-CHAIN` | **`SPEC-DRAFTER-CHAIN` cites llama.cpp by bare line number with no revision beside it.** Two sites landed by W1 at `31cefe631`: `include/vllm/config/speculative.h:34` cites `common/arg.cpp:3754-3763`, and `.agents/specs/drafter-chain.md` `## Upstream chain` cites that plus `common/speculative.cpp:2164-2186`. `.agents/oracles/llama-cpp.md` records `pin = 10bf611e533d81f739128304991c5e133c6aebd8`, `pin_label = b10451`, `pinned_on = 2026-08-16`, `gateable = no`, and neither citation names it. llama.cpp moves several times a day, so those ranges will name different code within weeks with nothing in the tree able to notice: `scripts/check-symbol-anchors.py` states outright that it cannot verify a LINE citation, the same gap that let `SPEC-DSPARK-QWEN3-ROUTING` carry a wrongly shifted anchor under a helper's name until a repair wave caught it by hand. **The exposure is bounded and stated**: W1 cites NO llama.cpp gate input — its rules are this engine's own document-shape decisions — so this is design context going stale, not a measurement resting on a moving target, which is why it is filed rather than fixed inside a wave scoped to a config field. Minimum repair: append `@ 10bf611e` to both citations and prefer a symbol over a line range where one exists (`common_speculative_n_max` in the same paragraph already needs nothing). Listed under `## Owed` in [drafter-chain.md](../specs/drafter-chain.md). Found by a fresh review of `row/SPEC-DRAFTER-CHAIN-W1` | gap | -| [#1608](https://github.com/mudler/vllm.cpp/issues/1608) | `BACKEND-ROCM` | **`main` is RED on a plain CPU build at its own tip.** Re-derived on a PRISTINE `origin/main` worktree at `e2a9e035d`, `cmake -DVLLM_CPP_BUILD_TESTS=ON`, no patches: `tests/vllm/v1/worker/test_runner.cpp:1557` `CHECK_THROWS_WITH_AS(make_runner(), "Block size must be a multiple of 16", std::invalid_argument)` receives `No valid attention backend for device type 0 from {CPU_ATTN: [block_size not supported], FLASH_ATTN: [block_size not supported]} (use_mla=false, use_sparse=false)` instead — 20 cases / 544 assertions / 1 failed, `Status: FAILURE!`. Deterministic, not a `ctest -j` flake: reproduced serially on two independent build trees with identical output. `e2a9e035d` (`test(BACKEND-ROCM)` [#1273](https://github.com/mudler/vllm.cpp/pull/1273), `#41`) added exactly those 29 lines and touched no other file, and its body records validation on gfx1151 / ROCm 7.2.3 at `20/20 cases, 544 assertions` — the count matches this run exactly, so the case ran and passed THERE. It cannot pass on CPU because `validate_configuration` (`BACKEND-ATTN-REGISTRY` [#1332](https://github.com/mudler/vllm.cpp/issues/1332), landed `369ea7fd4`) rejects both CPU-visible backends on `block_size not supported` BEFORE `initialize_kv_cache` reaches the multiple-of-16 check the case names. The guarantee is enforced on ROCm and a DIFFERENT guard gets there first on CPU — the "state was not the one you believed" shape, where the validation was real and the host was not the one CI runs. Two candidate repairs, neither chosen: accept either message (which weakens the case to "something refused" and stops distinguishing the two guards, the one thing it exists to do), or assert the registry message on a CPU build and the block-size message where a backend accepts the block size. The second keeps both guarantees pinned and is the larger edit; either changes what a landed case asserts, so it wants its own row, a red-before reproduction and green-after evidence. NOT fixed in flow by the `SPEC-DRAFTER-CHAIN` W1 repair that found it, which owns a `--speculative-config` field and touches nothing under `src/vllm/v1/worker/` or `src/vllm/attention/`. Found by merging `origin/main` into `row/SPEC-DRAFTER-CHAIN-W1`, running the full gate on the merged tree, then re-deriving on pristine `main` to establish it was not the merge | bug | -| [#1597](https://github.com/mudler/vllm.cpp/issues/1597) | `QUANT-QWEN38-27B-NVFP4-ARM` | **The Qwen3.5 dense safetensors loader picks each projection's quantized arm from which tensor NAMES are present, and never from what the checkpoint DECLARES.** `qwen3_5_dense_weights.cpp` routes with `IsNvfp4Projection` (`has(.weight_packed) \|\| has(.weight_scale_2)`), then `IsFp8BlockProjection`, then an `F8_E4M3` dtype probe, then bf16. W5 (#821) made the ModelOpt `MIXED_PRECISION` declaration readable at load and CROSS-CHECKS it, refusing a disagreement by name; it deliberately did not make the declaration authoritative, because doing so moves an arm on a gate model. Two consequences are live rather than hypothetical, both read from the hub rather than argued. **(1)** `VT_MODELOPT_W4A4=1` consumes `.input_scale` whenever the tensor exists, flipping `IsTrueW4A4()` to the fp4-ACTIVATION GEMM. `nvidia/Qwen3.6-27B-NVFP4`@`0893e1606ff3d5f97a441f405d5fc541a6bdf404` declares all 193 of its NVFP4 modules `W4A16_NVFP4` and ships an `input_scale` on every one, so the lever selects W4A4 against the declaration on the FP8-tower gate model of #466 — the checkpoint every recorded 27B NVFP4 ratio was taken on, which makes a fix a measurement change as well as a correctness one. `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121`@`36f717a22990e82c54c1d48ee77c491b87825680` declares the same algorithm and ships NO `input_scale` on an NVFP4 projection, so there the same lever silently does nothing and says nothing. **(2)** `IsQwen27QuantizedLinear` returns false for every `.linear_attn.in_proj_*` and for `lm_head`: right for `unsloth/Qwen3.6-27B-NVFP4`, whose `ignore` list covers `in_proj_{qkv,z,a,b}`; wrong for both ModelOpt 27B artifacts, which declare `in_proj_qkv` and `in_proj_z` FP8; wrong again for the r0b0tlab one, whose `lm_head` is `W4A16_NVFP4`. It is exported through `include/vllm/model_executor/models/qwen3_5_dense.h` and has NO production call site — only `tests/vllm/models/test_qwen27_dense_forward.cpp` — so it states a per-artifact fact as a per-model one and nothing loads through it. A fix routes by the declaration where one exists, keeps the name probe for checkpoints that declare nothing, and owes a same-binary A/B on `nvidia/Qwen3.6-27B-NVFP4` before and after, because the point of the change is that some module's arm moves. Listed under `## Owed` in [`qwen38-27b-quant-arms.md`](../specs/qwen38-27b-quant-arms.md) | bug | -| [#1602](https://github.com/mudler/vllm.cpp/issues/1602) | `BACKEND-ROCM` | **`e2a9e035d`, the current tip of `main`, added one `test_runner` case that is RED on a default CPU build.** `tests/vllm/v1/worker/test_runner.cpp:1557` asserts `CHECK_THROWS_WITH_AS(make_runner(), "Block size must be a multiple of 16", std::invalid_argument)`; with `block_size = 8` and no ROCm and no FLASH_ATTN the runner never reaches `CheckKvCacheShape`, because the attention-backend SELECTOR rejects every candidate first and raises `No valid attention backend for device type 0 from {CPU_ATTN: [block_size not supported], FLASH_ATTN: [block_size not supported]}`. Measured at `e2a9e035d` on a Release CPU build: the file is 20 cases / 544 assertions, 19 pass, this one fails, and over a 585-test `ctest` run it is the ONLY `***Failed`. The commit body records its validation as "on gfx1151 (Strix Halo, ROCm 7.2.3)", which is the one configuration where the selector has a backend that reaches the shape check, so the case encodes a build configuration it never declares. Both refusals are correct behaviour and they are DIFFERENT guarantees, so the fix is to assert whichever refusal the built configuration produces, or to gate the case on a build that has a backend supporting the block size and SKIP LOUDLY otherwise — never to widen the assertion to accept any throw. Found by `QUANT-QWEN38-27B-NVFP4-ARM` W5 ([#821](https://github.com/mudler/vllm.cpp/issues/821)) running the full suite for an unrelated loader change, in a file that change does not touch. Related: [#41](https://github.com/mudler/vllm.cpp/issues/41), [#1332](https://github.com/mudler/vllm.cpp/issues/1332) | bug | -| [#1593](https://github.com/mudler/vllm.cpp/issues/1593) | `KV-FP8` | **`KV-FP8` W2 and W3: the CUDA fp8 KV store, its paged-attention read, and the runner integration.** W1 landed the CPU half (`vt::ReshapeAndCacheFp8`, the read dequant in CPU paged attention, `vllm::v1::ParseCacheDType`) and left W2/W3/W4 `later`. The issue is now the critical path of benchmark campaign [#1574](https://github.com/mudler/vllm.cpp/issues/1574), whose subject `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` declares `kv_cache_quant_algo: "FP8"` and carries ZERO `k_scale`/`v_scale` tensors, so every published profile serves it with `--kv-cache-dtype fp8` and no cell can be served correctly without this. **W2 IS LANDED HERE**: the CUDA fp8-e4m3 store (`src/vt/cuda/cuda_cache.cu`), the fp8 dequant on the CUDA paged-attention read (`src/vt/cuda/cuda_paged_attn.cu` `LoadKv` + `LaunchPagedFp8`), the removal of the two W1 device-class refusals that made the CUDA arm unreachable however well it was registered, and a named CPU-or-CUDA refusal for the READ because it rides ADDITIVE `PagedAttentionArgs` fields on an op `kMETAL`/`kROCM` already register for the FLOAT path — without which an fp8 cache would be read as that backend's float dtype and return silent garbage. Gate `tests/vt/test_cuda_fp8_kv_cache.cpp`, RED-first on the provider-routing case. **The device half of that gate is UNEXECUTED and the CUDA TUs are UNCOMPILED**: the implementing session had no `nvcc` and no device, and says so under `## Owed` in [fp8-kv-cache.md](../specs/fp8-kv-cache.md) together with the reachability debt — nothing calls the fp8 KV path from a production entry point on either backend, which is **W3's** wiring (half-sized KV blocks, `--kv-cache-dtype` threading, the checkpoint scale path including this checkpoint's scales-absent case). W3, W4, the Metal/ROCm arms and fp8_e5m2 remain owed | feature | -| [#1636](https://github.com/mudler/vllm.cpp/issues/1636) | `KV-FP8` | **`KV-FP8` W1's three read-side comments anchor `scaled_vec_conversion` at `quant_utils.cuh:302-308`, which at pin `555967922` is the IDENTITY primary template plus the header of the fp8->HALF specialization.** Lines 301-305 are `template ... { return x; }` and 307-314 are the `` conversion; the `` one the comments describe is at `:419-429` under the `// fp8 -> float` label at `:418`. Sites, all landed by W1 and all outside the W2 change's authority: `include/vt/fp8_kv.h:92`, `include/vt/ops.h:1129`, `src/vt/cpu/cpu_paged_attn.cpp:164`. W2 ([#1593](https://github.com/mudler/vllm.cpp/issues/1593), PR [#1606](https://github.com/mudler/vllm.cpp/pull/1606)) copied the same wrong anchor into four new places and CORRECTED all four there; these three are filed rather than fixed in flow. Same shape, second anchor: `Fp8KVCacheDataType` is cited at `dtype_fp8.cuh:9-13`, which is the `#include ` guard -- the enum is at `:15-19` (`include/vt/fp8_kv.h:5`, `:30`). Third, a different kind: `.agents/engine-matrix.md` and `.agents/quantization-matrix.md` both say the W2 CUDA translation units are UNCOMPILED, and CI job `cuda-fat-build` built them for ten architectures under `-Werror=all-warnings` and PASSED on `4d71e776efc18cb5e61a26e642ddad8de5339134` (run 32495320287, job 96812232428). What stays true is that nothing has been EXECUTED on a device, because that job configures `-DVLLM_CPP_BUILD_TESTS=OFF`; both clauses need the narrower statement. An upstream anchor is how the next reader checks a port against the oracle, and one that lands on a `return x;` primary template invites the conclusion that the port is unfaithful. Listed under `## Owed` in [fp8-kv-cache.md](../specs/fp8-kv-cache.md) | bug | -| [#1594](https://github.com/mudler/vllm.cpp/issues/1594) | `BACKEND-BENCH-CUDA-SGLANG-PREFLIGHT` | **`run_serve_low.py` could not name a third benchmark subject, and the flag that would name one is the key of the evidence tree.** `--model-key` was `choices=("27", "35")` at `:663`, and the value is a directory name at `:75` (`evidence/corpus//`), `:81` (`evidence/raw///`) and again at `:350`/`:352` for the in-container paths the pinned `sglang.bench_serving` reads and writes. So [#1574](https://github.com/mudler/vllm.cpp/issues/1574)'s subject, `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` @ `36f717a2`, could only be measured by mislabelling its evidence under an existing key. **The failure to prevent is not a crash**: it is a run that completes, writes a plausible raw result, and files it under another subject's key, which a `choices` edit alone would not catch. FIXED IN FLOW. `MODEL_KEYS` is now one definition the parser's `choices` and a new fail-closed `build_bench_command` refusal both read, and the third key is `q38mtp` -- deliberately not digit-prefixed, because `27` already names a DIFFERENT 27B checkpoint (`unsloth/Qwen3.6-27B-NVFP4`) and `online_gate.py`'s `27`/`27n` pair records that two 27B checkpoints share no goldens and no comparable ratios. The test enters through `main()` on a real argv with a stubbed `subprocess.run` that resolves the container `--output-file` back through the command's OWN bind mount, so the host derivation and the container derivation must agree or the run cannot read its own result back; it asserts, for every admitted key, that all four derived paths carry that key and that no other admitted key's tree was named or created. Three negative mutations recorded red and restored: the corpus path hardcoded to `27` (`missing corpus partition`), the container raw path alone hardcoded to `27` (host read `FileNotFoundError` -- the insidious one the corpus precondition cannot see), and the key guard deleted. No workload shape, concurrency ladder or timed semantic changed, and no server was started. `docs/USAGE.md` is deliberately unedited: it documents no `tools/bench` harness, this flag reaches no shipped command, and adding the first `tools/bench` entry for one choice value would be a new obligation rather than a projection of a changed one | gap | -| [#1605](https://github.com/mudler/vllm.cpp/issues/1605) | `BACKEND-BENCH-CUDA-SGLANG-PREFLIGHT` | **`test_serve_low_client.py`'s concurrency-cap case reds at random, and the number it reports belongs to the fixture rather than to the client.** `tests/tools/test_serve_low_client.py:135` asserts `_CompletionHandler.peak <= 2` after `run_usage_batch(..., max_concurrency=2)`. Measured at `f7d517c09`: **1 failure in 30** standalone runs of that one case, `AssertionError: 3 not less than or equal to 2`; the fresh review of [#1596](https://github.com/mudler/vllm.cpp/pull/1596) measured 2 in 8 on the same head. `tests/CMakeLists.txt:12` registers all of `tests/tools` as the CTest target `test_serve_low_tools`, so it reds CI at random on whatever diff is in flight, and it has done so since `e58858a91` with no issue naming it. The cap itself is structural and is NOT what fails: `run_usage_batch` submits into a `ThreadPoolExecutor(max_workers=max_concurrency)`. `peak` counts handler threads inside the fixture's `do_POST`, between `active += 1` and the `finally: active -= 1`, and that window is strictly WIDER than the client's -- `post_json` returns once `response.read()` has the body and closes the connection in its own `finally`, while the handler thread has yet to leave its `try` and retake the lock. A worker that returns and submits its next request in that gap makes a third handler live with two client threads. So the reported 3 is the instrument's own overlap, and reporting it as a breached concurrency cap points the reader at `run_usage_batch`, where the 3 did not come from. **NOT FIXED IN FLOW, deliberately**: the repair changes what the case measures (count the client's in-flight window, or hold the handler open until the client has read), which changes the guarantee it asserts, so under `## Every change starts from an issue` it takes the row, spec and fresh-review path rather than riding inside a model-key change. The lower bound `assertGreaterEqual(peak, 2)` is the half that proves parallelism happened and must survive any repair: widening the bound to `<= max_concurrency + 1` would make the case green and measure nothing. Found by the fresh review of [#1596](https://github.com/mudler/vllm.cpp/pull/1596) (F4), reproduced independently before filing | bug | -| [#1595](https://github.com/mudler/vllm.cpp/issues/1595) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | `tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547` calls `::getpid()` with no `` in its include block, so the TU does not compile under clang-20 (`no member named 'getpid' in the global namespace`); landed on main in the SPEC-DFLASH2 GGUF arm and blocks every post-merge rebuild on a clang host, including the TT card gates; fixed in-flow with the one-line missing include, matching the sibling-test convention | bug | -| [#1576](https://github.com/mudler/vllm.cpp/issues/1576) | `LTX25-DEVICE-RESIDENCY` | **`test_ltx2_video`'s `span_slack` bound is a 250 µs wall-clock budget with an allowance for sanitizers and NONE for CPU load, and the asserted quantity swung 171x between two runs of ONE binary.** NOT CURRENTLY RED, and recorded before it becomes an intermittent nobody can attribute: `build-test-cpu` is green on `main` in run [32464632413](https://github.com/mudler/vllm.cpp/actions/runs/32464632413) at `483cd3198`, `100% tests passed, 0 tests failed out of 584`, with `test_ltx2_video` Passed in 201.56 s. `6b48edb2c` (GATE-CI-RED-REPAIR, [#1499](https://github.com/mudler/vllm.cpp/issues/1499)/[#1494](https://github.com/mudler/vllm.cpp/issues/1494)) repaired the suite and INTRODUCED the bound in the same change: `span_bound = min(kSpanSlackPerRecord * leaf_records, 0.5 * leaf_seconds)` with `kSpanSlackPerRecord` 250 µs (`tests/vllm/multimodal/test_ltx2_video.cpp:3827`, `:3830`), raised to 3 ms under either sanitizer (`:3822`, `:3825`). Measured at `origin/main` `6b48edb2c` on a detached worktree, CPU-only, x86_64, `BUILD_RC=0`, SAME binary both samples: in the full 102-case `ctest -j 1` run at load 82-110 the `artifacts.frames` leaf reported **3.223 ms** of slack on a 4.935 ms leaf (65.3%) against the 250 µs bound and the CHECK failed 13x over; the same case alone under `-tc=` at load ~90 reported **18.8 µs** on a 2.98 ms leaf (0.63%) and passed 13x under. The suite read `102 cases, 101 passed, 1 failed`, `4194 assertions, 5 failed`, every failure inside *"the three carrying phases contain their work and the load keeps its order"* -- `span_slack <= span_bound` (`:3909`) and `covered >= c.min_coverage * leaf_seconds` (`:3945`). Same commit, same tree: `test_ltx2_video` runs 201.56 s on the CI runner and took 2863.35 s on that box, 14x slower. The bound is a FIXED absolute time charged per instrument boundary against a leaf that can be milliseconds long, so whether it holds is a property of SCHEDULER LATENCY, and a descheduled thread is off-CPU for far more than 250 µs; the sanitizer arm already concedes the quantity is environment-sensitive and contention moves it the same way with no allowance. Cost is attribution, not the red: a reviewer who sees this fail cannot separate a genuine phase-instrumentation regression from a loaded runner, which is [#433](https://github.com/mudler/vllm.cpp/issues/433)'s class. NOT FIXED IN FLOW, deliberately: a bigger constant only moves the threshold at which the same ambiguity appears, and the candidate repairs (assert the containment RELATION rather than a duration, derive the budget from measured per-boundary cost on the running host, or gate the case on a quiescence precondition and report "not measured" instead of "failed") are a numerics/instrument decision owned by the LTX-2.5 residency lane rather than by the row that found it. Found while measuring [#1464](https://github.com/mudler/vllm.cpp/issues/1464) at `origin/main` | bug | -| [#1385](https://github.com/mudler/vllm.cpp/issues/1385) | `GATE-CI-AARCH64-COVERAGE` | **CI builds 4 of 552 test targets on aarch64, and every fleet GPU box and two release bundles are aarch64.** Re-derived at `e2a9e035d`: `tests/CMakeLists.txt` defines **552** `vllm_cpp_add_test` targets plus 32 direct `add_test(NAME ...)` registrations, and the x86-64 lane ran **584** CTest entries with 3 skipped on run 32465485947. `build-test-cpu-arm64` (`.github/workflows/ci.yml:1096`) builds four ISA and kernel-tier targets and runs no `ctest` at all, so 0.72 % of the suite executes on the architecture the project measures, gates and ships on. Its own flag block (`:1110`) additionally sets `VLLM_CPP_BUILD_EXAMPLES=OFF` and `VLLM_CPP_SERVER=OFF`, which puts `examples/tokenize` and the `/v1/completions` parse segment out of reach of any target-list change. Decided in [ci-aarch64-coverage.md](../specs/ci-aarch64-coverage.md): a curated subset is REJECTED because a stated sensitivity principle (weak memory model, `char` signedness, floating-point contraction, hash order) selects **291 of 552** targets, so half a suite costs most of a full build; a self-hosted fleet runner is REJECTED because it takes a leased box outside `rc`. Recommended: one new `build-test-cpu-arm64-full` job that builds everything and runs `ctest` serially, landing `schedule`+`workflow_dispatch` only with `continue-on-error`, then promoted per-PR once measured. Measured cost: **$0** (`timing` reports `total_ms: 0`; the repository is public), **+45-55 job-minutes** against a median of 357 per scheduled run, and **zero** added wall-clock while it finishes inside `cuda-fat-build`'s measured 123.0-minute median finish. It makes the hermetic tokenizer parity goldens execute on aarch64 for the first time, which is the `## Owed` item [prompt-token-divergence.md](../specs/prompt-token-divergence.md) names, but the committed corpus carries only **30** combining marks over 99 lines against the 74-150 per prompt that produced the anomaly, so it is a necessary and not a sufficient probe | bug | -| [#1617](https://github.com/mudler/vllm.cpp/issues/1617) | `SPEC-DFLASH2` | **`d27639e71` re-added an unguarded `#include ` to `tests/vllm/models/test_qwen3_dflash2_gguf.cpp`, a file [#1581](https://github.com/mudler/vllm.cpp/pull/1581) (`13548db8f`) had already fixed through the `process_id` seam.** The include supports NOTHING: the file carries no `::getpid` call and includes `support/process_id.h` at line 68. It is the [#603](https://github.com/mudler/vllm.cpp/issues/603) shape of the [#503](https://github.com/mudler/vllm.cpp/issues/503) class -- MSVC ships no ``, so an unguarded include does not fail on Windows, it does not COMPILE, and `tests/support/process_id.h` says so in its own header comment. It did not surface as a new red because both `windows-msvc-*` lanes are baseline-red and never run on `main`. ORDERING is the whole cause: the change was authored against the pre-#1581 tree, where it was correct, and landed after #1581 had removed the call it existed to support. FIXED IN FLOW by deleting the one line; found while merging external contributor pull requests, not owned by that work | bug | -| [#1613](https://github.com/mudler/vllm.cpp/issues/1613) | `GATE-QWEN38-27B-FP8-BLOCK` | **The `Qwen/Qwen3.8-27B-FP8` block-wise token gate cannot be taken, because the 28.75 GiB checkpoint is not on the share.** `/mnt/nas_share/rc/ckpt/` holds `qwen3.8-27b-hf`, which is the **bf16** artifact -- no `quantization_config` key, `text_config.dtype = bfloat16` -- and `qwen3.8-q1_0`. Neither is this subject. The share has 3.4 TiB free, so the cost is AUTHORITY: `.agents/developer-preferences.md` authorizes large downloads for the `SPEC-DFLASH2` assets only. Nothing else blocks the gate, and that was not known before: a range-request audit of all 66 shard headers at revision `017b9c7a` shows every one of the 407 `F8_E4M3` tensors has `N % 128 == 0` and `K % 128 == 0`, so the sm120 complete-scale-block refusal (#1453) that makes DSV3's `kv_a_proj_with_mqa` unservable blocks NOTHING here; the ragged GDN `in_proj_a`/`in_proj_b` `[48, 5120]` are `BF16` and named in `modules_to_not_convert`; `weight_scale_inv` ships `BF16` (byte-checked via `data_offsets`, not the label) which `LoadFp8BlockRaw` already widens by value; and the per-layer `layers-.safetensors` naming already resolves through `SelectWeightFiles`. Spec `.agents/specs/gate-qwen38-27b-fp8-block.md`, parent #1189 | gap | -| [#1614](https://github.com/mudler/vllm.cpp/issues/1614) | `GATE-QWEN38-27B-FP8-BLOCK` | **Three sites said `Qwen/Qwen3.8-27B-FP8` ships "~400" `modules_to_not_convert` entries, and at revision `017b9c7a` it ships 882** (882 unique, 636 outside the vision tower). The number is the evidence for an ARGUMENT -- it is why `IsFp8BlockProjection` reads the config AND the tensors instead of probing dtypes -- so being wrong by more than 2.2x invites the next reader to re-derive it. No reading of the list produces ~400: the visual entries are duplicated under two naming conventions, so distinct modules are about 759, and half of 882 is 441. Sites: the comment above `IsFp8BlockProjection`, the comment above `Fp8BlockQuantConfig::modules_to_not_convert`, and `.agents/specs/model-fp8-block-weight.md`. The routing itself is correct and no defect in it is asserted; two other claims in the same comment were checked against the checkpoint headers and hold (zero `input_scale` tensors, and the `[96, 40]` block-grid hazard is real). Found while auditing the checkpoint for #1613, fixed in the same flow | bug | -| [#1502](https://github.com/mudler/vllm.cpp/issues/1502) | `VT-REFTIER-HOST-ADDRESSABLE` | **`docs/ENVIRONMENT.md` described `VT_ADOPT_DEVICE_BYTES` as Vulkan-only and said it has "No effect on CUDA/CPU/Metal", and [`cffe59b02`](https://github.com/mudler/vllm.cpp/commit/cffe59b02) ([#1477](https://github.com/mudler/vllm.cpp/issues/1477)) made both halves false.** That change moved `ReferenceTierEligible` off `UnifiedMemory()` onto `Backend::DeviceMemoryIsHostAddressable()` and added truthful overrides so no backend lost the reference tier, so `MetalBackend` now answers `MetalContext::unified_memory()` and `RocmBackend` answers its `unified_memory_`. The weight loader gates the lever on exactly that predicate, at both `AdoptDeviceBytesAsHost` branches in `src/vllm/model_executor/models/qwen3_5_weights.cpp`, so the lever ACTS on Apple silicon and on an integrated ROCm part. **The correction is not "add two backend names".** Every number in that row is GB10 through Vulkan, and nobody has measured the lever on either new arm, so the row now separates the backends it is MEASURED on from the backends that merely satisfy the predicate — reach and measurement are different claims and the row read as if the measurement covered the reach. CUDA and CPU stay inert and are unchanged: neither overrides the default `false`, which `tests/vllm/platforms/test_platform.cpp` pins for GB10, and the CPU backend answering `UnifiedMemory() == true` while the narrower predicate stays `false` is the whole reason the two properties are separate. The MEASUREMENT on Metal and integrated ROCm stays owed and is listed under `## Owed` in [`vt-reference-tier-host-addressable.md`](../specs/vt-reference-tier-host-addressable.md); it needs an Apple-silicon box or an integrated AMD part | documentation | -| [#1628](https://github.com/mudler/vllm.cpp/issues/1628) | `SPEC-DFLASH2` | **The DFlash2 candidate selector could not consume a QUANTIZED target `lm_head`, so the arm was refused on the one checkpoint the `BENCH-QWEN38-27B-SOTA` campaign has.** Measured on `dgx:gpu0` 2026-08-21: this engine loads `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` @ `36f717a2` and generates correctly (canary `19 x 23 -> 437`, warm decode 11.06 tok/s vs vLLM's 9.71 on the same box), and attaching the DFlash2 draft died at `vllm_engine_load: dflash: target tensor lm_head.weight is not BF16 (got U8)`. The draft owns no head — it runs the TARGET's — and `SharedHeadSource` read it with one `LoadNamedBf16`, refusing on the STORED DTYPE. That predicate cannot separate the two states `## Risks/decisions` D12 is about: a head WIDENED into something the target does not compute with, and a head kept PACKED and computed with natively. D12 is NOT reversed — it stands for the GGUF container, the only one that still widens a head — and the safetensors arm now takes the target loader's OWN routing decision (`DenseLmHeadTakesNvfp4`) so an NVFP4 head lands packed in `Qwen3DFlashWeights::lm_head_fp4` and the draft's logits GEMM is the same W4A16 dispatcher the target's head takes. The merged oracle agrees and needs no branch: at vllm-project/vllm#52816 head `b389ac29` `compute_candidates` carries no quant-method check and goes through `LogitsProcessor.get_top_k_tokens` -> `_apply_head` -> `lm_head.quant_method.apply`, so the guard this port mirrored at head `66e5414c` is gone. Gated on the PROPERTY and not the load: the draft's block forward must be BITWISE equal to `Qwen3_5MTPModel::ComputeLogits` — the other draft that shares the target's head — over the same hidden states, so the selector's top-K is the target's exactly. FP8 and true-W4A4 heads still refuse by name, DSpark still refuses (`## Owed` O28), and the CUDA arm plus the real checkpoint are owed a measurement (`## Owed` O29) | bug | -| [#1635](https://github.com/mudler/vllm.cpp/issues/1635) | `VT-REFTIER-HOST-ADDRESSABLE` | **[#1620](https://github.com/mudler/vllm.cpp/pull/1620) cited `tests/vllm/platforms/test_platform.cpp` as pinning the CUDA backend's `DeviceMemoryIsHostAddressable()` default, and that test reads the CPU backend.** `FakeUnifiedAddressablePlatform`'s `device_type()` override returns `DeviceType::kCUDA` while its `backend()` override returns `vt::GetBackend(DeviceType::kCPU)`, so the assertion `CHECK_FALSE(gb10.backend().DeviceMemoryIsHostAddressable())` reads CPU and the test's own comment says so. Grep those three names; do not look for a line number, because this index is append-only and a number written HERE is permanent whatever file it points at. **Nothing in the tree pins the real `CudaBackend`.** The CONCLUSION survives — `CudaBackend` declares no override, so it inherits `Backend::DeviceMemoryIsHostAddressable`'s `false` default in `include/vt/backend.h` — but it holds by absence, which is a weaker claim than a pin and must not read as one. No gate could catch this: the assertion passes, and it passes for a reason unrelated to what it claims to measure. The wrong citation landed in three places and one is PERMANENT: the `VT_ADOPT_DEVICE_BYTES` row of `docs/ENVIRONMENT.md` (corrected in this flow), the commit message, and the `#1502` row of this index, which is append-only and can never be edited — which is why this correction needs a row of its own rather than an edit. Both are named by ROW and not by line on purpose: the first draft of this row cited the `#1502` row by LINE NUMBER, and `origin/main` appended rows above it while this branch was open, so the line that draft named had stopped identifying the `#1502` row before this row had shipped even once, and a permanent line number into an append-only file would have been wrong on the day it landed — which is [#844](https://github.com/mudler/vllm.cpp/issues/844)'s rule, committed inside the row that exists to describe this class of defect. No COUNT of the rows that moved it is written here either: that number is read off whatever `origin/main` happened to be merged last, so it rots by the very mechanism this row condemns, and a draft of this row carried one that was already stale by the next merge. Found by post-hoc review of four merged records pull requests, not by a gate. The pull request that lands this row carries NO closing keyword and #1635 stays OPEN: this change corrects the citation and records the gap, while the pin itself needs a CUDA device. OWED: either pin the real `CudaBackend` or state in the record that the default holds unpinned | bug | -| [#1629](https://github.com/mudler/vllm.cpp/issues/1629) | `KERNEL-ATTN-DENSE-FLASH` | **`test_check_attention_rung_consistency.py` stored a count of the model tree, so every row on the attention-rung allowlist redded it by doing the thing the allowlist exists for.** `ShippedTreeTests::test_the_population_is_not_empty` asserted `>= 9` against a tree holding exactly 9 `vt::Attention(` sites, so a removing row had zero headroom and no green path: leaving the parked stem redded the floor (`8 not greater than or equal to 9`), and deleting it redded the floor and `test_allowlist_holds_only_the_in_flight_stems` as well -- while the allowlist header explicitly recommends the first of those two. That is the `## Records` shape AGENTS.md names, a measurement of one file stored inside another, and it blocked PR #1579 (#1545) and the LTX-2.5 routing row, which removes two of the three parked stems. FIXED HERE, and NOT by lowering the number, which is the known mute-switch: the floor became `>= 1`, because an empty population means the scanner broke and that is the only thing a raw total can honestly detect, and the guard that a rename cannot slip past stays `test_the_six_deliberate_sites_carry_a_marker`, which pins six sites BY NAME. A case that the stem in a red message names a real source file was added beside it, so a typo in the allowlist is still caught without pinning a count. A second drift lock in the same suite, `assertGreater(excused, 0)`, required the shipped allowlist to stay non-empty forever; it is replaced by two synthetic cases that build their own allowlisted population, so the excused counter is pinned without the shipped tree having to keep a stem parked. Found while landing #1578 and #1579 together -- each green in isolation, main red once both land -- and fixed in the same flow | bug | -| [#1631](https://github.com/mudler/vllm.cpp/issues/1631) | — | **A comment-only edit is impossible in any of the 43 `scripts/check-*.py` checkers, so a comment that is measurably false in one cannot be corrected.** `scripts/check-pr-size.py:170` classifies every `scripts/check-*.py` and `scripts/check-*.sh` as a `governance_checker`, and `change_errors` then demands a paired `tests/scripts/test_*.py` change that `executable_evidence` proves goes RED against the BASE checker. A comment-only diff leaves BASE and HEAD semantically identical, so no test can distinguish them and no such evidence can exist. Measured on this row: `ERROR: BASE checker stayed green for 'scripts/check-attention-rung-consistency.py'; changed test is not semantic evidence`, rc=1, with the identical invocation against the parent commit exiting 0. Live cost, three comments in `scripts/check-attention-rung-consistency.py` that ship unrepaired in #1578: `:58-61` says widening to `\bAttention\s*\(` is not the repair because it would match every fast rung, when the reason a wider pattern is not the repair is the function-pointer call it still cannot reach; `:93-96` says the `\b` is what excludes `vt::AttentionDenseFlash(`, when the trailing `\(` is, and the `\b` only excludes a leading identifier character as in `xyvt::Attention(`; `:252-255` says `sites - marked` is not the excused count, when on this tree it is (9 sites, 6 marked, 3 excused). The suite beside them was repaired for #1629, so the tree now contradicts itself across two files in the same directory pair. NOT fixed in the flow that filed it: teaching the guard to tell a comment-only or docstring-only diff from a semantic one changes what the gate accepts, which AGENTS.md `## Changing the rules or a checker` routes to its own row, spec and red-before evidence, and the honest report is therefore a filed gap rather than a comment smuggled in beside an unrelated semantic change. A candidate patch is parked on the issue, and two smaller pre-existing defects in `check-pr-size.py` itself (an incomplete entry-point list at `:370-371`, an unread `SELF_CHECKER` constant at `:376`, which the issue body records as `:378` because a line anchor drifts inside the pull request that writes it) are frozen by the same lock. Owed under `## Owed` in [attention-rung-visibility.md](../specs/attention-rung-visibility.md) | bug | -| [#1632](https://github.com/mudler/vllm.cpp/issues/1632) | `QUANT-QWEN38-27B-NVFP4-ARM` | **W6's NVFP4 token gate named [#1185](https://github.com/mudler/vllm.cpp/issues/1185) as the authority it waits on, and #1185 closed on 2026-08-18 as local-only** -- it tracked one operator's machines rather than a defect here -- so five sites pointed a reader at an issue that reports "closed" without reporting "cleared": `docs/FEATURES.md`, and the spec's `**Related:**` header, wave table, blockers section, `## Owed` list and `## Now`. **The blocker did not close with the issue, and it is not the one the citations described.** The pinned oracle `5559679229bc961848b121ccdeaa8fa5d79bec98` DOES build, install, import and GENERATE TOKENS inside an `rc` lease on `dgx:gpu0` (2026-08-18), which kills the "a model run is untested" clause those sites carried, and #1213 killed the "a lease cannot produce a runtime" premise underneath it. It survived at `max_num_batched_tokens` 512, `max_model_len` 512 and `gpu_memory_utilization` 0.30 on a ~20 GiB model, where the recorded denominator for this family is 8192 and 2048; `AGENTS.md` §Gates requires vLLM's PRODUCTION configuration as the denominator, so a reduced-`mnbt` arm is a different engine setup rather than a smaller measurement, and `gpu_memory_utilization` is a REFUTED lever (`.agents/specs/mtp-k-gt-1.md`: 0.75 thrashed 42 minutes, 0.30 rebooted the box). The named next levers are `max_num_batched_tokens` and `cudagraph_capture_sizes`, one at a time. The second half is the bytes: `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121`@`36f717a2` is ~20.4 GiB over four shards and is not mirrored where a lease can read it, which is also why its sha256 is recorded as unpaid. Same shape as [#1613](https://github.com/mudler/vllm.cpp/issues/1613) for the block-wise FP8 gate. FIXED IN FLOW: all five citations now name this issue, and the loader is untouched -- W5's accounting and cross-check need no lease and no oracle. Spec [`qwen38-27b-quant-arms.md`](../specs/qwen38-27b-quant-arms.md), parent [#821](https://github.com/mudler/vllm.cpp/issues/821) | gap | -| [#1538](https://github.com/mudler/vllm.cpp/issues/1538) | `SPEC-DFLASH2` | **vllm#52816's head moved a THIRD time (`66e5414c` -> `3406ec1d`) while it is still open, and refactors `compute_candidates` into `LogitsProcessor.get_top_k_tokens`.** Measured 2026-08-21 by W6 from `raw.githubusercontent.com` at both heads: +11/-80 on `qwen3_dflash2.py`, +4/-16 on `dflash2/speculator.py`, +2/-5 on the base `speculator.py`. The big one is a RELOCATION rather than new math -- the padding mask, the id rebase, the TP all-gather and the scale-THEN-softcap order all survive in `logits_processor.py:241-286`, so `## Owed` O16's reading of the codebook-span question holds at BOTH heads. What does NOT survive is the explicit `UnquantizedEmbeddingMethod`/`UnquantizedLinearMethod` guard that `## Risks/decisions` D12 ports as `RefuseQuantizedDflash2LmHead`, which is deleted at `3406ec1d`; our guard's own reason (the GGUF arm dequantizes `output.weight` to bf16, and a GGUF target with a safetensors DFlash2 draft is admitted here) is independent of upstream's and stands. NOT reconciled in flow, deliberately: `## Gates` G2 fixes the gate head at `66e5414c` while the pull request is unmerged, and moving the port onto a third unmerged head during the gate would move the thing being measured. Owed under `## Owed` O21 of [the DFlash2 spec](../specs/dflash2-spec-decode.md) | verification | -| [#1456](https://github.com/mudler/vllm.cpp/issues/1456) | `SPEC-DFLASH2` | **The GB10 oracle DOES have a FLASH_ATTN denominator: the arch measurement stands, the conclusion drawn from it does not.** Measured 2026-08-21 by W6 on `dgx:gpu0` through an `rc` lease, with the very wheel #1456 was filed about (`vllm-0.1.dev1+g66e5414c6`, sha256 `fbc247ab...`). A capture that exported `VLLM_ATTENTION_BACKEND=TRITON_ATTN` got `FLASH_ATTN` anyway and RAN: `Using FlashAttention version 2`, 54.87 GiB loaded, CUDA graphs captured (PIECEWISE 5/5, FULL 1/1, plus the DFlash2 speculator's own), 4 x 64 coherent tokens, speculation live at 209 accepted of 350 drafted and mean acceptance length 5.00. No `cudaErrorUnsupportedPtxVersion`. Consistent with the `sm_80`/`sm_75` SASS finding rather than contradicting it: `sm_80` PTX JITs FORWARD, and that error is the OPPOSITE failure (PTX newer than the driver). So `FA_USABLE=0` in the staged `FA-CONSTRAINT.txt` was inferred from emitted arches, never observed from a run, and is the thing to reconcile. A SECOND trap found in the same run and recorded so nobody repeats it: **`VLLM_ATTENTION_BACKEND` does not exist at this revision** -- grepping every `.py` in the wheel returns nothing; the knob is `EngineArgs.attention_backend` (`arg_utils.py:706`) folded into `AttentionConfig.backend` (`:2382`), so the old export selects NOTHING and auto-selection wins silently, letting a run record one backend while executing another. NOT reconciled in flow: W6 does not substitute a denominator the developer declared, and takes both arms instead, each named in its own golden. Owed under `## Owed` O22 of [the DFlash2 spec](../specs/dflash2-spec-decode.md) | verification | -| [#1561](https://github.com/mudler/vllm.cpp/issues/1561) | `SPEC-DFLASH2` | **vllm#52816 MERGED, so this row's gate head and O21's parked D12 decision both come due.** Read from the forge 2026-08-21: `merged: true`, `merged_at 2026-08-21T05:27:22Z`, merge commit `b389ac29465b33f9e9c534df221ea3c129e9793f`, head `3406ec1dae9916f920b90f0dbf90dcf54923d042`. The merge landed 46 minutes BEFORE `SPEC-DFLASH2` W6's work commit `bb416e0ae` was authored (`06:13:50Z`), so five statements in the spec plus the `#1538` index row recorded an open pull request that had already closed. `## Gates` G2's own rule is "`66e5414c` if #52816 has not merged, and the merge commit if it has", so the head it selects today is `b389ac29`. W6's capture stays pinned to `66e5414c` because that is the wheel that ran and it predates the merge -- a dated exception, not the rule -- and what is owed here is moving the gate head and re-reading G2 and G3 at the merged head. O21 additionally parked the `UnquantizedEmbeddingMethod`/`UnquantizedLinearMethod` guard deletion as "the decision when #52816 settles"; it has settled onto vLLM's `main`, our `RefuseQuantizedDflash2LmHead` stands on its own independent reason (the GGUF arm dequantizes `output.weight` to bf16), and writing that decision down against a merged upstream rather than a branch is owed. Re-verified at `3406ec1d`: `logits_processor.py:241-286` preserves the padding mask, the id rebase, the TP all-gather and the scale-THEN-softcap order, so O16's reading holds at the merged head too. Records corrected on `row/SPEC-DFLASH2-W6`; the WORK is owed here | verification | -| [#1562](https://github.com/mudler/vllm.cpp/issues/1562) | `SPEC-DFLASH2` | **W6's oracle capture harness exists only as PROSE, and the FLASH_ATTN label is a post-hoc relabel of an uncommitted log.** `## Owed` O22 and O23 were written so the next agent would not pay three 51.75 GiB loads again; they describe the hook on `DFlashSpeculator.propose` below the `cg_mode == FULL` branch, the `torch.cuda.is_current_stream_capturing()` delegation, the resolved-backend read-back and the abort-on-zero -- and none of it is in the tree, nor is `w6-relabel.py`, nor any run log. Three consequences visible in the committed evidence: `dflash2_27b_spec_on_flash_attn.json` carries `attention_backend_source: "corrected from the run log by w6-relabel.py; the capture's original value came from VLLM_ATTENTION_BACKEND, which does not exist in this wheel and selected nothing"`, which does not meet the read-back rule O22 itself lays down and cannot be re-derived; the TRITON_ATTN golden's `hook_stats` reads `{propose_calls: 59, skipped_dummy: 1, skipped_capture: 0}` against 55 recorded blocks with contiguous `call` ids 3..57, so `59-1-0 = 58 != 55` and THREE propose calls are unaccounted for (now bounded one-sidedly by the gate and pinned at 3, still unexplained); and `## Owed` O17's discharge has no committed runner and no log, so its peak-RSS pair and its `[SPECTRACE]` lines cannot be re-derived either. Owed: commit the harness, the relabel script and the log, or record that they were lost with the lease | bug | -| [#1564](https://github.com/mudler/vllm.cpp/issues/1564) | `SPEC-DFLASH2` | **The two divergent draft blocks were attributed to the SELECTOR's rank contraction, and the block shape argues against it.** W6 measured 45 of 47 draft blocks byte-identical and wrote that both flips are "the lattice op is a REDUCTION over `selector_rank`". Nothing measured that: the golden records `{call, req_row, anchor, drafts}` per block and no values, no logits and no top-2 gap, so it cannot say whether either flip was a near-tie at all, let alone in which reduction. The shape points the other way -- in BOTH blocks only slot 2 changes while slots 3-6 are byte-identical, and `src/vt/cpu/cpu_ops.cpp:3219` has step l read block row `previous`, the slot step l-1 chose, so a flipped CHILD INDEX would move the predecessor row every later step reads and four identical later slots would be four coincidences per block, twice. A different candidate ID at the SAME winning slot, a rank swap in `ComputeCandidates`' top-k over the target head's logits, produces this shape with none. `SPEC-DFLASH` D6 licenses a near-tie envelope; it does not license labelling an unmeasured flip as one nor naming the op. The attribution is WITHDRAWN in the spec and the benchmark record rather than replaced. The instrument is available and cheap: `Qwen3DFlash2Model::ComputeCandidates` already returns `(ids, values)` and so does upstream's `compute_candidates`, so the next capture records the top-2 candidate margin at the flipping slot on both sides. Blocks nothing -- the gate reads 45/47 on a majority bar and both blocks emitted the same target tokens either way | bug | -| [#1607](https://github.com/mudler/vllm.cpp/issues/1607) | `SPEC-DFLASH2` | **The DFlash2 startup notice prints TWICE on every draft load, and `docs/USAGE.md` called it one-time.** `CheckDflash2DraftArm` (`src/vllm/entrypoints/model_loader.cpp:502`) ends in an unconditional `std::cerr <<` of the whole notice paragraph with no once-flag, and the loader reaches it TWICE on one load of one `EngineParams`: directly from `FromModelDir` at `:1929`, deliberately placed ahead of every path, config, tokenizer and weight operation so a misclassified draft is caught before a 51.75 GiB target is mapped; and again from `ResolveSpecConfig` at `:1206`, which the `LoadedEngine` constructor runs in its member initializer at `:1538` on all three `new LoadedEngine(...)` returns (`:2172`, `:2341`, `:2359`). The server, the C ABI and the bench client therefore each emit the paragraph twice, on the safetensors arm and the GGUF arm alike. The tree already states that the resolution re-runs (`model_loader.cpp:2313`, `:871-874`); what nothing stated is that the notice re-runs with it. ESTABLISHED STATICALLY, by reading the call graph rather than by executing it -- a runtime confirmation needs a DFlash2 checkpoint and a rebuild of the whole 464-object library. Cosmetic rather than behavioural: nothing is loaded twice and nothing is refused twice, and `CheckDflash2DraftArm` returns early for every non-DFlash2 draft. Found by the THIRD fresh review of `SPEC-DFLASH2` W6 (#1314) and NOT fixed in flow, deliberately: the fix changes the production loader and needs its own red-first test and its own fresh review, while that wave was scoped to prose. `docs/USAGE.md` is corrected in the same change, so the shipped documentation is no longer wrong about the behaviour while this is open. Owed under `## Owed` O25 of [the DFlash2 spec](../specs/dflash2-spec-decode.md) | bug | -| [#1654](https://github.com/mudler/vllm.cpp/issues/1654) | `ENG-LOAD-DIRECT-UPLOAD` | **`WeightF32` (`src/vllm/model_executor/models/qwen3_5.cpp:1054-1060`) forms a `const uint16_t*` over an `OwnedTensor`'s bytes, which are a BORROW of the safetensors mapping whenever `BorrowStTensorBytes` took the direct-upload arm.** A safetensors data offset carries no 2-byte alignment guarantee, so the read is undefined behaviour. This is the [#301](https://github.com/mudler/vllm.cpp/issues/301) / [#772](https://github.com/mudler/vllm.cpp/issues/772) class at a site neither reached: both were closed by repairing the sites they had FOUND rather than by giving the borrow a contract, and #772's own title says one of its four was "invisible to UBSan by construction". MEASURED, not read: [#1633](https://github.com/mudler/vllm.cpp/pull/1633)'s loader-seam cases drive `LoadedEngine::FromModelDir` against a synthetic on-disk Qwen3.5 dense target, and while their fixture writer left `8 + header_len` odd the `sanitize-cpu (address,undefined)` lane reported `qwen3_5.cpp:1059:84: runtime error: load of misaligned address ... for type 'const short unsigned int', which requires 2 byte alignment`; that lane builds `-fno-sanitize-recover=all` (`CMakeLists.txt:260`), so the process exits 1 at the first finding. `BorrowStTensorBytes` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:431-461`) already PROMISES what is missing -- it says it will "FAIL CLOSED on anything that is not a whole-range verbatim view" and that "a mismatch can only cost the lever, never correctness" -- and checks the keep-alive, the data pointer, the byte count and the exact span, but never ALIGNMENT, the one mismatch that costs correctness. NOT fixed in flow because the two candidate repairs trade differently and the choice is a load-time measurement this row owns: refusing the borrow on an unaligned source honours the stated contract but silently disables the direct-upload lever for a whole checkpoint, while reading through `vt::LoadUnaligned` keeps the lever and leaves the next such site to be found the same way. Reachability from a PUBLISHED checkpoint is not claimed: the reference `safetensors` writer pads its header, and nothing here measures what every producer on the hub does; what is measured is that the loader accepts an unpadded producer and then reads it as UB. #1633 pads its own fixture header instead, and says so at the writer. Spec [`load-direct-upload.md`](../specs/load-direct-upload.md) | bug | -| [#1646](https://github.com/mudler/vllm.cpp/issues/1646) | `SPEC-DFLASH2` | **`tests/tools/` held 351 cases across 20 suites and NO lane ran one of them.** Measured at `e100e64e1` on a clean worktree: `python3 -m unittest discover -s tests/tools -t . -p "test_*.py"` reports `Ran 351 tests in 21.067s / OK`, standard library only, zero skips, no GPU and no vLLM wheel — and no workflow, no CTest registration and no `scripts/agent-preflight.sh` line executed it. The only `unittest` invocations in `.github/workflows/` are four `tests.scripts.*` modules; preflight's `SUITES` loop runs `tests/scripts/$suite.py` only; a tree-wide grep for `tests.tools` outside the directory returns prose in `.agents/` and four unrelated path constants in `scripts/check-snapshot-pins.py`. Worse than untested: the suites are QUOTED AS EVIDENCE — `.agents/parity-ledger.md` carries "all tools 34/34" on five `SERVE-GATE-ONLINE` rows and `.agents/upstream-sync.md:38` records "34 of the 233 `tests/tools` cases" — so they read as gating in every document that cites them while being reachable only by an agent who typed the command. What they cover is not marginal: `test_oracle_pin.py` is the [#520](https://github.com/mudler/vllm.cpp/issues/520) oracle-identity assertion, `test_gpu_clock_state.py` is the [#543](https://github.com/mudler/vllm.cpp/issues/543) clock attribution every ratio rests on. FIXED IN FLOW by one preflight line and one CI step, DISCOVERED rather than enumerated because an enumerated list is a shared file every new suite must edit, which is the record-lock shape `AGENTS.md` §Records forbids. Found while wiring [#1562](https://github.com/mudler/vllm.cpp/issues/1562)'s refusal gate, which would otherwise have landed dead; not owned by that row's subject matter | bug | -| [#1648](https://github.com/mudler/vllm.cpp/issues/1648) | `SPEC-DFLASH2` | **[#1646](https://github.com/mudler/vllm.cpp/issues/1646)'s central claim is FALSE: `tests/tools/` HAS a CTest registration and CI has been running it on every pull request.** #1646's row says the 351 cases across 20 suites had "no workflow, no CTest registration and no `scripts/agent-preflight.sh` line" executing them. `tests/CMakeLists.txt:12-16` registers the target `test_serve_low_tools` running exactly `python3 -m unittest discover -s tests/tools -p "test_*.py"` with `PYTHONPATH` set, and has since `e58858a91`; `CMakeLists.txt:2728` makes it live and `.github/workflows/ci.yml:1057` runs `ctest --test-dir build` inside `build-test-cpu` on every `pull_request`. Executing that command on the branch head reports `Ran 414 tests / OK` in 20.2 s. The error came from grepping `tests.tools`, the DOTTED module path, while CMake and the workflow spell it `tests/tools` -- a null grep proving the terms wrong rather than the thing absent. Three consequences: `.agents/parity-ledger.md`'s `all tools` citations and `.agents/upstream-sync.md:38` were citing a LIVE suite, not a dead one; the NARROW gap was real and its fix stands, because preflight genuinely ran none of them, and the CI step's remaining value is a Python-only verdict in ~21 s under its own name rather than behind a C++ build; and the DISCOVERY mechanism #1646 chose is correct and unaffected. TWO FURTHER ERRORS in the same immutable row, recorded here because it cannot be edited: "pinned at 3" pins nothing, since `hook_reasons` asserts `recorded_blocks <= propose_calls - skipped`, which holds for any residual at or above 0; and "all tools 34/34 on five rows" does not occur -- the ledger reads `35/35` on two rows and `164/164` on eleven, and 34 is `upstream-sync.md`'s count of cases that go RED under a released-wheel version shape, a different quantity from a pass count. The #1646 row is append-only and STANDS AS WRITTEN; this issue and `## Owed` O27 of [the DFlash2 spec](../specs/dflash2-spec-decode.md) are the authority over it, as this campaign handled the #1538 row. Prose corrected in `.github/workflows/ci.yml` and `scripts/agent-preflight.sh` in the same change. Found by the fresh review of the SPEC-DFLASH2 speed harness ([#1562](https://github.com/mudler/vllm.cpp/issues/1562)) | bug | -| [#1657](https://github.com/mudler/vllm.cpp/issues/1657) | `SPEC-DFLASH2` | **The SPEC-DFLASH2 speed gate can NEVER emit a number: the clock window is written only AFTER the arm it must describe.** Measured on `dgx:gpu0` 2026-08-22 over leases `11cee02a`/`52ac5673`/`a03f34e4`, running the committed gate at `bed3feae6` unmodified. `open_clock_window` started `tools/bench/gpu_clock_state.py sample` and handed the arm `--clock-summary `; that module writes the summary only when the sampler STOPS, which is after the arm, so `_load_clock` raised `cannot read the clock summary .../clock-vllm.json: [Errno 2] No such file or directory` BEFORE the model loaded. Pre-closing the window does not help either, because `clock_reasons` floors a window at 30 busy samples and 50% busy and a pre-run window is idle: `every one of 98 clock samples was idle; there is no window to attribute the measurement to`. The summary had to describe the arm AND exist before it, and both cannot hold. **No test caught it** because `tests/tools/test_dflash2_speed_harness.py` pre-wrote `clock.json` in `setUp`, so every case received a summary that already existed and nothing drove the DRIVER's ordering -- the state was not the one anybody believed, a green suite over an instrument that cannot run. FIXED IN FLOW: the clock is a precondition of the MEASUREMENT and not of the arm's execution, so the ARM owns its window (`dflash2_oracle_capture.ClockWindow`, which runs `gpu_clock_state`'s own `sample` CLI) -- it runs, its sampler stops, the summary is written, and only then is it read and judged. Two consequences are deliberate: the arm record is written BEFORE the verdict, because a leased arm costs about two hours and discarding its evidence makes the next run pay the same lease, and nothing quotable is emitted since the refusal precedes the `print` and `build_speed_result` refuses the record again through `clock_pairing`; and the oracle arm's window opens AFTER `LLM(...)` has loaded, so the 12-minute load is outside the samples that must be 50% busy. Our arm's window still spans four `vllm-cli` loads, because that binary is one process per prompt -- recorded, not worked around. The suite now pre-writes nothing, a stub sampler writes the summary on STOP, and a `.running` marker makes "is the window open right now" observable, because the samples file survives the sampler and reading it would answer yes to a window that had already closed. Owned by [`dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) `## Owed` O28 | bug | -| [#1658](https://github.com/mudler/vllm.cpp/issues/1658) | `SPEC-DFLASH2` | **`BACKEND_PROBES` misses the beyond-pin wheel, and one scalar backend under-describes this model -- three resolve at once.** Measured on `dgx:gpu0` 2026-08-22 from inside the unmodified committed arm, after a 54.87 GiB weight load, a 42.7 GiB KV allocation and CUDA-graph capture. All three committed walks raised `AttributeError: 'GPUModelRunner' object has no attribute 'attn_backend'`; the refusal named every probe, took no fallback and invented no label, which is the designed behaviour and DISCHARGES `## Owed` O26 residual 1. The repair is one entry, measured against the live engine: `llm_engine.vllm_config.attention_config.backend` and `llm_engine.engine_core.engine_core.vllm_config.attention_config.backend` each return `TRITON_ATTN`. `self.attn_backend` no longer exists on either runner in this wheel; it lives on the `Attention` layer. **The larger point:** `...worker.model_runner.attn_groups` resolves THREE backends at once on this model -- `GDNAttentionBackend` over the 30 `linear_attn` layers, `TritonAttentionBackend` over the 16 full-attention layers, and `FlashAttentionBackend` over the DFlash2 draft's five sliding-window layers (`model.layers.64-68`) -- and the denominator's backend is exactly what W6 measured vLLM disagreeing with ITSELF over (0.597 against 0.657 acceptance). FIXED IN FLOW: the two measured walks lead `BACKEND_PROBES` and the three retired ones are KEPT after them, because the list is ordered, an older wheel still answers on them and a walk that resolves nothing costs one `AttributeError`; and the per-group map is RECORDED beside the scalar while only the SCALAR is gated -- the scalar is what the run declared, what `attention_backend_reasons` compares against and what `test_qwen38_dflash2_spec_decode.cpp` reads off a golden, and the two spellings are not comparable (`TRITON_ATTN` against `TritonAttentionBackend`), so a checker equating them would refuse every correct run. A missed group walk is a named `miss`, never an empty map, because an empty map reads as one backend over every layer. STILL OPEN: the run watched the draft's five layers resolve `FlashAttentionBackend` and generate tokens while `FA-CONSTRAINT.txt` records `FA_USABLE=0` on sm_12x from [#1456](https://github.com/mudler/vllm.cpp/issues/1456), so #1456 needs a re-read before that constraint is quoted again. Owned by [`dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) `## Owed` O28 | bug | -| [#1659](https://github.com/mudler/vllm.cpp/issues/1659) | `SPEC-DFLASH2` | **The speed harness asserts `--attention-backend` and never passes it to `LLM()`, so the declared denominator is unreachable by any path.** Measured on `dgx:gpu0` 2026-08-22 at `bed3feae6`. `capture()` built `LLM(...)` with no backend kwarg while `attention_backend_reasons` requires `resolved == declared`, and the arm log under a declared `TRITON_ATTN` reads `Using FLASH_ATTN attention backend out of potential backends: ['FLASH_ATTN','FLASHINFER','TRITON_ATTN','FLEX_ATTENTION']`. So repairing [#1658](https://github.com/mudler/vllm.cpp/issues/1658) ALONE would make the arm resolve `FLASH_ATTN`, compare it against the declared `TRITON_ATTN` and refuse on the mismatch: a working gate needs this, the probe entry and [#1657](https://github.com/mudler/vllm.cpp/issues/1657), not any one of them. This reproduces W6's recorded failure at a later head -- [#1456](https://github.com/mudler/vllm.cpp/issues/1456) measured vLLM's vendored flash-attention unable to target sm_12x, which is why `TRITON_ATTN` is the declared oracle backend on this box at all. FIXED IN FLOW: the declared backend is passed over the spellings `ATTENTION_BACKEND_KWARGS` names -- `attention_config` carrying a `backend` key first, because the measured read-back walk is `vllm_config.attention_config.backend`, then a bare `attention_backend`. The spelling is UNVERIFIED at the beyond-pin head, on the same footing `BACKEND_PROBES` was on: a wheel REJECTS a kwarg it does not declare with a `TypeError` raised while `EngineArgs` is built and therefore before anything loads, so trying both costs no lease time; a wheel that takes NEITHER is a loud refusal naming both; a spelling that is accepted and IGNORED is caught by the read-back, which is this refusal; and `--attention-backend-kwarg` pins the answer once known, with no code change. Owned by [`dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) `## Owed` O28 | bug | -| [#1660](https://github.com/mudler/vllm.cpp/issues/1660) | `SPEC-DFLASH2` | **A DFlash2 lease needs `cuda-libraries-dev-13-0` and `python3-dev`, and `RC_LEASE_ID` does not exist on this fleet.** Measured on `dgx:gpu0` 2026-08-22 over leases `52ac5673` and `a03f34e4`. DFlash2's `compute_candidates` -> `_topk` -> `flashinfer.topk` JIT-compiles `topk.cu`, which includes ``; the `cuda-toolkit-13-0` metapackage does NOT install that header, and leg B died on it INSIDE `profile_run` **after a 12-minute model load**, presenting as a model failure rather than a missing header. Leg C installed `cuda-libraries-dev-13-0` and got past it. `python3-dev` is likewise required or Triton driver JIT compilation fails, and that failure surfaces as `Model architectures ['Qwen3_5ForConditionalGeneration'] failed to be inspected` -- which names the model, not the toolchain. Separately, the leased worker carries `RC_DEVICE`, `RC_JOB_ID` and `RC_TOKEN` and no `RC_LEASE_ID`, so the gate's `--lease-id` default read empty and the `## Owed` O26 recipe run VERBATIM refused with `no lease id. Claim the device with rc run/rc hold first; never ssh to a fleet box`. The refusal is correct on a missing lease; the defect is that the committed procedure could not satisfy its own gate. FIXED IN FLOW: the default takes `RC_JOB_ID` first and falls back to `RC_LEASE_ID` -- the measured variable first, the documented-but-absent one for a controller that does export it -- both packages are recorded in [`environment.md`](../environment.md) beside the existing note that no CUDA toolkit is preinstalled, and the O26 recipe now names them. Owned by [`dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) `## Owed` O28 | bug | -| [#1543](https://github.com/mudler/vllm.cpp/issues/1543) | `ENG-RECORD-CONFLICT-SURFACES` | `check-public-doc-tables.py` gates the projection documents for size and shape only, never for key uniqueness, so `docs/STATUS.md`, `docs/BENCHMARKS.md` and `docs/FEATURES.md` can each carry two contradictory rows for one key with nothing indicating which is current. MEASURED on [PR #1081](https://github.com/mudler/vllm.cpp/pull/1081): `docs/BENCHMARKS.md` carried the key `**NemotronH paged forward** (MODEL-NEMOTRON-H-ABI-A2P, #810)` TWICE and the two rows contradict -- one says the GB10 read was 4/24 with the sm_121a re-run pending a lease, the other says the same gate PASSES 96/96 -- and the checker returned rc=0 on that file. Counted directly, the key appears once on `origin/main` and twice on the PR head. The blindness is PROVED BY MUTATION rather than inferred from reading: a reviewer reinserted a duplicate key with contradictory content and `benchmarks_errors()` still returned `[]`, and the checker holds no uniqueness or duplicate-key logic at all (positive control: 14 `errors.append` calls in the same file, so the grep pattern finds what is there). No checker asserts uniqueness ON A PROJECTION DOCUMENT, and that is the whole gap -- duplicate-key checks themselves are neither scarce nor novel here, and the closest is prior art the fixing row should mirror rather than reinvent. `scripts/check-agent-record.py::check_matrices` already refuses a `duplicate ID` over the matrix inventories, keyed on the row ID and reported against the file and line that repeats it, and `scripts/check-agent-record.py::check_issue_index` refuses `issue #N listed twice` -- the refusal this very row had to be checked against before it was appended. That same file READS `docs/BENCHMARKS.md`, through `scripts/check-agent-record.py::markdown_files`, but never for key uniqueness. The remaining duplicate-key checks are structural rather than record-keyed and transfer nothing: a duplicate YAML key in `check-release-workflow.py`, a duplicate JSON key in `ready-for-helper.py`, a duplicated oracle-pin key in `check-oracle-pins.py`. The cause is that two record surfaces carry OPPOSITE disciplines and nothing tells an author which one is open: `.agents/issue-index.md` is an append-only log where appending is correct, and `docs/BENCHMARKS.md` is a keyed projection where appending is the defect. Two review passes and one operator verification missed it for one reason worth recording -- each checked that both sides' hunks survived the merge, which proves the merge was faithful and says nothing about whether the RESULT is key-unique. SCOPE for the fixing row, so it does not over-promise: assert uniqueness on the first table cell and scope it PER TABLE, because two tables may legitimately share a key and because `audit-live-rows --check` already reports three IDs live in two matrices (BACKEND-CPU, BACKEND-CUDA-SM121, BACKEND-VULKAN), so a naive global check reds `main` on day one -- the failure this repository already suffered twice in one week in [#1464](https://github.com/mudler/vllm.cpp/issues/1464) and [#1468](https://github.com/mudler/vllm.cpp/issues/1468). The red-before is available and is not hypothetical: PR #1081's own pre-repair head contains a real duplicate. NOT fixed in flow, and this index row does not fix it -- a checker-semantics change owes its own row, spec, red-before test or mutation, and green-after evidence, and #1081's separate repair collapses its two rows into one WITHOUT closing this, because the next append reproduces it. Owned here because this row is the standing owner of public-document projection-gate semantics ([#460](https://github.com/mudler/vllm.cpp/issues/460), [#495](https://github.com/mudler/vllm.cpp/issues/495), [#498](https://github.com/mudler/vllm.cpp/issues/498), [#507](https://github.com/mudler/vllm.cpp/issues/507), [#1055](https://github.com/mudler/vllm.cpp/issues/1055)) -- named at that width deliberately, because the five do NOT all sit in one script: #460's `check_links` and #498's retired `MAX_README_CHARS` live in `check-agent-record.py` and `check-readme-structure.py`, so calling the group `check-public-doc-tables.py` semantics would be false of two of them. The row is `READY` rather than `DONE`, and its spec already names `tests/scripts/test_check_public_doc_tables.py` as a red-before suite, so the owner is a live row whose declared test surface is the checker this issue indicts; the fixing row still owes the spec and the red-before, which ownership records rather than replaces | bug | -| [#1663](https://github.com/mudler/vllm.cpp/issues/1663) | `KERNEL-ATTN-DENSE-FLASH` | **Three stems outlived their removing rows on `scripts/attention-rung-allowlist.txt`, and while a stem sits there it excuses its WHOLE translation unit -- including a marker somebody deletes.** `47a918d8f` ([#1579](https://github.com/mudler/vllm.cpp/pull/1579), issue [#1545](https://github.com/mudler/vllm.cpp/issues/1545)) routed `muse_glimmer_vision`'s perception encoder to `vt::AttentionDenseFlash`, and `90e8c3c85` ([#1557](https://github.com/mudler/vllm.cpp/pull/1557), issue [#1549](https://github.com/mudler/vllm.cpp/issues/1549)) swapped the LTX-2.5 DiT device forward to the same op; after both, `muse_glimmer_vision` names `vt::Attention` nowhere, and `ltx2`/`ltx2_device` name it only at calls that now carry their own `VT-ATTN-NAIVE:` reason (the host CPU-only arm, where `kAttention` and `kAttentionDenseFlash` are the same registered function, and the `VLLM_LTX2_DIT_FLASH_ATTN=0` arm of a same-binary A/B). `scripts/check-attention-rung-consistency.py` reported all three `STALE (not a failure)` and exited 0, which is the header's design -- the removing row is not forced to edit this file and the deletion falls to the next preflight -- so the entries were owed rather than wrong. **The hole is MEASURED, not inferred.** On `db648fb88`, deleting the `// VT-ATTN-NAIVE:` marker at `src/vllm/model_executor/models/ltx2.cpp:959` leaves the checker at **rc=0** (`7 carry a recorded reason, 1 unmarked and excused`) with the three stems listed, and at **rc=1** naming `ltx2.cpp:966` with them removed; both arms restored byte-for-byte against a pre-taken sha256. The window was open from `90e8c3c85` until this fix. FIXED IN FLOW: the stems are deleted and `tests/scripts/test_check_attention_rung_consistency.py::ShippedTreeTests::test_allowlist_holds_only_the_in_flight_stems` moves to the empty set in the same change, as the allowlist header requires. The empty set is not a weaker assertion -- `drift_sites` now excuses nothing, so `test_shipped_tree_is_green` measures the tree on its markers alone -- and a new case asserts the three stems positively and ASYMMETRICALLY, `ltx2`/`ltx2_device` present-and-marked against `muse_glimmer_vision` absent, because one assertion covering all three is false of one of them in either direction. Distinct from [#1629](https://github.com/mudler/vllm.cpp/issues/1629), which is the population floor in the same checker and is untouched here. Spec [`attention-rung-visibility.md`](../specs/attention-rung-visibility.md), parent [#1544](https://github.com/mudler/vllm.cpp/issues/1544) | bug | -| [#1470](https://github.com/mudler/vllm.cpp/issues/1470) | `LTX25-PHASE-RESIDUE` | **`test_ltx2_video` false-redded once on `main` under load and the failing case's identity was never captured.** The prediction it records (F12 of #1441's fifth review) was that a coverage floor whose per-boundary sampling cost is fixed while the leaf shrinks is a false-RED risk and never a false pass. Measured by `LTX25-PHASE-RESIDUE`: the prediction is right about the polarity and WRONG about the term. The `denoise` coverage miss is not sampling cost, it is the sampler's post-process and Euler step, which no anchor wrapped and which scale with the latent -- 49 us per step at nine frames against 343 us at 81, in ONE run of one binary, and instrument cost does not move 7x with the latent. **STILL OPEN, and the correction above is the whole of what this row establishes.** The row's branch anchored that work as `denoise.update` and [#1556](https://github.com/mudler/vllm.cpp/issues/1556) is CLOSED rather than merged, so nothing anchors it on `main` today; the anchor is owed under [#1668](https://github.com/mudler/vllm.cpp/issues/1668). An earlier draft of this row said the issue was closed by anchoring, which was true of the branch and never of the tree. Spec [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1536](https://github.com/mudler/vllm.cpp/issues/1536) | `LTX25-PHASE-RESIDUE` | **`test_ltx2_video` was persistently red rather than load-flaky, and it was the only failing test on `main`.** The issue asks for the residue to be DECOMPOSED rather than re-argued, and names `d995c52f0` (the temporal x2 upsampler) as the first hypothesis to test. The decomposition REFUTES that hypothesis and settles the cause: **92% of the un-named time is one region**, `Ltx2VideoEngine::Load` from the timeline's origin to `Open("load.dit")` -- 17.661 ms of a 19.178 ms residue -- while the upsampler's own work sits inside `phase.upsample_latent`, a named leaf that does not appear in the residue at all. The remaining four gaps are 4.95%, 1.30%, 1.09% and 0.56%, and the sixteen gaps between adjacent named phases hold 6.8 us each, which is the instrument and nothing else. **The RED is gone and the issue is NOT.** `519303d15` ([#1622](https://github.com/mudler/vllm.cpp/pull/1622), row `LTX25-DEVICE-RESIDENCY`) names that same 92% region `load.open` -- same open point, same close point, same `Scope::Close` shape as this row's `load.setup` -- so the sum floor is repaired on `main` by another row, and the coverage floor is repaired by `6b48edb2c` moving it to 0.75, which that change's own comment describes as a holding action in substance, though not in those words. What is owed is the naming that would make a tight floor honest again, under [#1668](https://github.com/mudler/vllm.cpp/issues/1668). Spec [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1567](https://github.com/mudler/vllm.cpp/issues/1567) | `LTX25-PHASE-RESIDUE` | the res_2s sampler arm has no `denoise.update` anchor, so its denoise would be decomposed on one arm only. Filed by `LTX25-PHASE-RESIDUE` and deliberately NOT landed with it, because no gate in this tree renders on that arm and an anchor no gate runs is dead code. `Ltx2Res2sDenoisingLoop` runs its own post-process and step behind `Ltx2Res2sHooks`, so the anchor needs a hook rather than a statement. It lives in `ltx2_samplers.cpp`, is declared in `ltx2_samplers.h` beside the hooks struct, and is called from `ltx2_video.cpp`. **NOT `ltx2_res2s.cpp`**: #1556's spec named that file and it has never existed here, which `git log --all --diff-filter=A` confirms; #1567's forge text names no file at all, so the wrong anchor came from the spec rather than from the issue. Doubly owed while [#1668](https://github.com/mudler/vllm.cpp/issues/1668) is open, since the first-order arm has no anchor either. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1568](https://github.com/mudler/vllm.cpp/issues/1568) | `LTX25-PHASE-RESIDUE` | `denoise.step` and `denoise.update` are open to a seconds transfer that no assertion in `test_ltx2_video.cpp` can see: (1b') compares `start_seconds` only, and no per-part floor separates the honest share (0.45% to 11.15% across four boxes) from a transfer (~0%). Leaving `denoise.step` open across the post-process and emitting `denoise.update` empty after it preserves the alternation, both counters, containment, non-overlap, exclusivity, (1c) and (2), while moving 100% of the decomposed seconds onto one name. Found by the fresh review of [#1536](https://github.com/mudler/vllm.cpp/issues/1536); `LTX25-PHASE-RESIDUE` claimed it closed and WITHDREW the claim when a reviewer measured it. Closing it needs an anchor INSIDE the callee. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1569](https://github.com/mudler/vllm.cpp/issues/1569) | `LTX25-PHASE-RESIDUE` | `PhaseLog::WriteJson`'s clock ordering is ungated and its own mutation stays GREEN 10 of 10 -- `wall 0.0608987s, unaccounted 0.000534223s, table charge 0.000301655s` -- because the copy and sort of a three-record table are nanoseconds. The claim was WITHDRAWN from the source by `LTX25-PHASE-RESIDUE` rather than defended, which is the reason this is a filed gap and not a passing test. Gating it needs a table with enough records for the sort to be measurable and a `WriteJson` with nothing between it and the last `Close`. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1570](https://github.com/mudler/vllm.cpp/issues/1570) | `LTX25-PHASE-RESIDUE` | nothing bounds the instrument's own share of a leaf, so an `uncovered <= 2 * leaf_instrument` bound can widen SILENTLY while printing a small number -- moving the DiT `Tick` out of `Evaluate` would charge ~110 flushed writes to `denoise` and buy a budget larger than the floor it replaces. Found as F5 by that row's fresh review. This is also what would close [#1439](https://github.com/mudler/vllm.cpp/issues/1439), whose filed complaint is that the budget is a SHARE OF `wall` and therefore decides by box load: what is needed is a bound on a quantity the scheduler cannot move, and `## Design` 3 of the spec is the measured evidence that the obvious candidate is not it. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1571](https://github.com/mudler/vllm.cpp/issues/1571) | `LTX25-PHASE-RESIDUE` | the phase table cannot be decomposed into its gaps without a scratch script, which is how [#1439](https://github.com/mudler/vllm.cpp/issues/1439) stayed open over a residue that turned out to be 92% ONE gap. NOT "three months", which #1556's spec and #1571's own forge title both say and which is impossible here: #1439 was filed 2026-08-20 and the repository's first commit is `accfae2de`, 2026-07-02. The four issues that argued the tolerance span under 24 hours. Filed by `LTX25-PHASE-RESIDUE`, which wrote that script and did not ship it, so a reader of `phase-log.json` still cannot see the decomposition and the same investigation will be re-derived the next time the residue moves. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1572](https://github.com/mudler/vllm.cpp/issues/1572) | `LTX25-PHASE-RESIDUE` | assertion (1c)'s span slack reds intermittently on `main` -- `decode.video` at `0.00256913` against a `0.00075` bound, 3.4x. Pre-existing from `6b48edb2c` and NOT `LTX25-PHASE-RESIDUE`'s, which does not touch (1c) and keeps its constants; observed by that row's fresh review and filed rather than repaired. Untouched by the record landing, since [#1556](https://github.com/mudler/vllm.cpp/issues/1556) is closed and no assertion moves. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1619](https://github.com/mudler/vllm.cpp/issues/1619) | `LTX25-PHASE-RESIDUE` | **The `merge=union` driver on this file DUPLICATES a row when both sides append before the same trailing anchor instead of at the true end, and no gate counts rows.** Measured on this row's merge of `origin/main` at `e2a9e035d`: base 521 lines, each side +8, correct union 537, driver produced **538** with `#1546` byte-identical at lines 523 and 533. `git merge-tree --write-tree` called it clean and `check-issue-index-append-only.py` passed, because a duplicate is an ADDITION and it only collects removals. **`check-agent-record.py` did NOT pass it**, contrary to what #1556's spec recorded: regenerating the driver output and running that tree's own checker returns rc=1 with `issue #1546 listed twice`, a refusal live since `8dd6508da` (2026-08-09). Reproduced twice, and restored byte-for-byte. The blind gate is ONE checker, not two. The repair side is narrower than #1556's spec claimed and is MEASURED here rather than reasoned: committing a duplicate leaves `check-issue-index-append-only.py` at rc=0, and committing its removal ALSO leaves it at rc=0, because that checker diffs `merge_base..HEAD` and an add-then-remove nets to nothing. The de-duplication half is CONDITIONAL, and the condition is what #1556's spec omitted: the checker reds a repair only when the DUPLICATE IS ALREADY IN THE BASE. Measured at three pairings -- `--base e2a9e035d` against the real canonical 537-line file rc=0, against a synthetic 537 rc=0, and `--base --head <537 de-dup>` rc=1. It diffs `merge-base..HEAD`, so when the base predates the duplicate the addition and the removal CANCEL and it passes. Since `origin/main` is preflight's base, and is the shape this branch used, the gate does NOT red someone who repairs driver output before committing it -- only someone repairing a corruption that already landed. The same range property is why relocating a base-reachable row DOES red it: moving row `#168` to the end gives rc=1 and a `removed:` line naming it. So 'de-duplicating in place fails the checker' is false unqualified, and true once the duplicate is base-reachable. GitHub does not honour the driver at all, so this bites local resolutions only and is invisible in pull request review. Canonical fix: take the merged branch's file byte-for-byte, append at the TRUE end, assert `git diff --numstat ` is `N 0`. Fix candidate is a multiplicity plus id-uniqueness checker, which fails 3 of 7 assertions on the raw driver output and names `#1546`. NOT shown to recur: #1556's spec said the same driver dropped `#838` on a later re-merge, and re-running `git merge-file --union` at every later merge where `#838` was on a side leaves it present in all of them, so that claim is WITHDRAWN rather than carried. A union driver cannot drop a line that is an addition on one side; a take-ours resolution can, which would indict the resolver and not the driver. Same record surface as [#1543](https://github.com/mudler/vllm.cpp/issues/1543), which is the projection documents' missing key-uniqueness gate; both are cases of a record surface whose checker cannot see a duplicate key. Owed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1668](https://github.com/mudler/vllm.cpp/issues/1668) | `LTX25-PHASE-RESIDUE` | **Three LTX-2.5 phase anchors and the phase instrument's own cost accounting were measured, gate-run and through three fresh reviews on [#1556](https://github.com/mudler/vllm.cpp/pull/1556), and NONE of them is on `main`.** Verified on `db648fb88` by a tree-wide grep over `src/`, `include/`, `tests/` and `docs/`: `load.dit_config`, `artifacts.mux`, `sampler_updates` and `Record::instrument_seconds` return ZERO occurrences, and `denoise.update` returns exactly one -- `tests/vllm/multimodal/test_ltx2_video.cpp:4325`, which is `6b48edb2c`'s own comment ASKING for it while calling its 0.75 floor a holding action. What IS on `main` is the load prologue, and only that: `519303d15` ([#1622](https://github.com/mudler/vllm.cpp/pull/1622), row `LTX25-DEVICE-RESIDENCY`) names it `load.open` at `src/vllm/multimodal/ltx2_video.cpp:794` and closes it at `:956`, which is the same region [#1556](https://github.com/mudler/vllm.cpp/pull/1556) called `load.setup`, opened and closed at the same two statements with the same `Scope::Close` shape. That supersedes ONE anchor, not the branch. The pull request is CLOSED rather than merged because `squash_merge_commit_message = PR_BODY` makes its body the permanent commit message, and that body argues at length for `load.setup` as new work -- merging it would have written a materially false narrative onto `main` irreversibly. Owed here: `load.dit_config` (0.950 ms, 4.95% of the measured 19.178 ms residue), `artifacts.mux` (0.210 ms, 1.09%), `denoise.update` plus `Ltx2ConditioningTrace::sampler_updates` (the WHOLE coverage miss, 49 us per step at nine frames against 343 us at 81 in one run of one binary), and `Record::instrument_seconds` with its conservation invariant. NOT urgent: neither red reproduces, the sum floor being correctly repaired by `519303d15`. DO NOT re-propose `residue <= 2 * instrument` -- three fresh reviews measured it red 4 in 45, 3 in 200, 2 in 200 and 28 in 160, and `## Design` 3 of the spec carries the distributions. The reference implementation, gate report and mutation table stay readable at `refs/pull/1556/head` = `b45ea3bbb`. Spec [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1671](https://github.com/mudler/vllm.cpp/issues/1671) | `SPEC-DFLASH2` | **Our arm's clock window was the PROCESS LIFETIME of four `vllm-cli` runs, so it was 18.37% busy and `clock_reasons` refused it.** Measured on `dgx:gpu0` 2026-08-22, lease `9ee9f53a`, at `04ed7b984`: `clock: ours was idle for 2630 of 3222 SM-clock samples (18.37% busy, below the 50% floor); the retained window does not describe the measured work`. The arithmetic is the arm's own legs -- window 3377 s, all 20 legs 1052.5 s, the 4 cold legs 959.3 s, the 16 warm legs **93.2 s**, and ~2325 s inside the window and inside no leg at all, because `vllm-cli` is one process per prompt and each reads a 52 GiB checkpoint off CIFS before it decodes anything. The window was ~3% warm generation. THE REFUSAL IS CORRECT AND THE FLOOR IS UNCHANGED: `MIN_BUSY_FRACTION` stays 0.5, `MIN_BUSY_SAMPLES` stays 30, and every threshold still applies -- to a smaller and truer set of samples. FIXED IN FLOW by moving the WINDOW rather than the rule. `examples/cli/main.cpp` prints a second line per leg, `run=%d/%d generate_start_unix=%.6f generate_end_unix=%.6f`, from `system_clock` (the timing line keeps `steady_clock`, because a duration and an instant are different questions and only the instant can be lined up against another process's samples). `gpu_clock_state` gains `sample_instant_s`, `normalize_spans`, `samples_within_spans`, `build_spanned_clock_record` and `read_sample_stream`, and `build_spanned_clock_record` DELEGATES to `build_clock_record` so the statistics, the idle accounting and the mid-window field check are the same code the unrestricted path runs. `dflash2_our_arm` gains `LEG_SPAN_RE`, `parse_leg_spans`, `legs_with_spans`, `warm_leg_spans` and `--clock-samples`, and `scripts/dflash2-speed-gate.sh` hands it the stream `open_clock_window ours` writes. The window is the WARM legs only, not all 20: `fold_legs` discards run 1 for a named cause, so a window that kept run 1's span would attribute the number to work the number excludes -- the same defect at a smaller scale. `is_warm_leg` is the one predicate both consumers read, so the number and the clock that qualifies it cannot drift onto different legs. Four refusals are ADDED, not relaxed: no span at all is refused rather than widened back to the whole stream, a reversed span is refused, a span that retained nothing is refused naming the stream it missed, and a leg printed with no marker refuses naming `examples/cli/main.cpp` so a binary built before the marker cannot silently drive this arm. The production `main.cpp` is LINKED against a stub libvllm and RUN in `CliMarkerRuntimeTest`, which asserts the printed instants land in wall-clock time -- the one check that catches a `steady_clock` epoch, which on Linux is BOOT and would make every span miss every sample. THIS DOES NOT MAKE THE GATE EMIT A NUMBER: the oracle arm still cannot read a summary that does not exist yet ([#1657](https://github.com/mudler/vllm.cpp/issues/1657)) and runs first, so the run still stops there. Discharges `## Owed` O30 item 2 of [the DFlash2 spec](../specs/dflash2-spec-decode.md) (authored as O28 item 2, which is what #1657's index rows say), recorded as O32 there | bug | -| [#1673](https://github.com/mudler/vllm.cpp/issues/1673) | `SPEC-DFLASH2` | **The two DFlash2 arms' clock windows now describe different leg populations, and `compare_clock_records` compares them anyway.** After [#1671](https://github.com/mudler/vllm.cpp/issues/1671) our arm's record covers the union of its 16 WARM leg spans and nothing else, while the oracle arm's covers whatever prompt loop the driver's window happened to span -- all 20 legs, the 4 it discards included, plus every gap between them. `median_offset_pct` and `mean_offset_pct` are both gated at 1.0%, and the mean is gated precisely because the excursion population does not cancel between the arms ([#1546](https://github.com/mudler/vllm.cpp/issues/1546)); two windows over different populations do not have that cancellation property at all, and the oracle's includes the cold legs on which its draft recorder is OPEN and pays two `.tolist()` copies per propose. On the 2026-08-22 run the oracle window read 85 samples at 98.82% busy so the difference was small -- a property of that run, not of the rule, and nothing measures it on the next one. NOT FIXED IN FLOW because it is not fixable on `main`: the oracle arm does not OWN its window, `scripts/dflash2-speed-gate.sh` opens the sampler before `LLM(...)` and closes it after the arm exits, and the arm never sees a finished stream, so there is no point at which a span restriction could be applied. #1657's `ClockWindow` repair (unmerged, `row/SPEC-DFLASH2-harness-fix`) is what creates that point. Closing it needs no new instrument on that side: `capture()` already times each leg, so it records the two instants from the same `time.time()` base the sampler stamps and calls the same `build_spanned_clock_record` and `warm_leg_spans` our arm calls, which is what makes the two arms one rule rather than two. Separately recorded and deliberately NOT filed: our arm pays four model loads to the oracle's one, which after #1671 is outside both the number and its clock -- `fold_legs` discards run 1 of each repetition group on BOTH arms, so both fold 16 warm legs of 20 -- leaving a second-order per-process difference (page cache, allocator, first touch) that nothing here sizes. `## Owed` O32 of [the DFlash2 spec](../specs/dflash2-spec-decode.md) | bug | -| [#1666](https://github.com/mudler/vllm.cpp/issues/1666) | `SPEC-DFLASH2` | **The O26C `attn_groups` census records 30 GDN layers; the run resolved 48, and 30 is the TEST STAND-IN's shape.** Four places -- `.agents/specs/dflash2-spec-decode.md` `## Owed` O30, the `BACKEND_GROUP_PROBES` comment in `tools/bench/dflash2_speed_harness.py`, `resolve_attention_backend_groups` in `tools/bench/dflash2_oracle_capture.py`, and `test_ONE_SCALAR_UNDER_DESCRIBES_this_model_so_the_MAP_is_recorded_too` -- recorded the 2026-08-22 leased walk as `GDNAttentionBackend` over **30** `linear_attn` layers. Re-derived from the run's own `c-probe-result.json`, under `candidate_walks[...worker.model_runner.attn_groups]`, by parsing every `AttentionGroup(backend=..., layer_names=[...])` pair and asserting the match count equals the 15 occurrences of `AttentionGroup(backend=`: **48 GDN layers in 10 groups** (eight of five, two of four), 16 `self_attn.attn` in 4 groups at every 4th index 3-63, and 5 draft layers (`model.layers.64-68`) in 1. The 16 and the 5 were right; only the 30 was wrong, and 30 is `_stand_in_attn_groups()`'s `range(30)` -- a fixture's arbitrary shape generalised into a measurement claim, on the row whose whole history is that failure class. The tree already disagreed with itself: `.agents/specs/qwen38-27b-quant-arms.md` and `tests/vllm/models/test_qwen38_27b_gguf_manifest.cpp` both say 48. Introduced by `04ed7b984`. FIXED IN FLOW: all four corrected to 48, and `_stand_in_attn_groups()` now carries the measured 10/4/1 group shape and the real layer names, because a fixture is what the next reader lifts a count from. That alignment closed a real hole rather than a cosmetic one -- turning the per-backend sum into an overwrite leaves the OLD one-group-per-backend fixture GREEN and reds the new one. A FIFTH copy is in the append-only #1658 row and cannot be edited; this row is its correction. Owned by [`dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) `## Owed` O30 | bug | -| [#1667](https://github.com/mudler/vllm.cpp/issues/1667) | `SPEC-DFLASH2` | **Both DFlash2 arms judge the LEGS before the arm record reaches disk, so one 63-token prompt discards ~2h of leased evidence.** `require_no_reasons(hook_reasons(...) + leg_reasons(...))` sits after the clock window and before the write: in `capture()` on the oracle arm, whose `main()` writes only after `capture()` returns, and in `main()` on ours. A non-clock refusal therefore takes `records`, `blocks`, `output_token_ids` and the O26 provenance with it. Proven by execution through the existing `OurArmRunEntryPointTest` fixture at `completion = 63`: `arm JSON on disk: ABSENT` against the control run's `PRESENT`. SEPARATE FROM #1657 and not a regression of it -- #1657 moved the CLOCK judgement after the write on exactly this reasoning and both arms carry the comment saying so, but that guarantee is scoped to the clock and says nothing about the leg and hook checks, while a leg refusal is the likelier one. `git log -S` places both call sites at `208559a79` (#1653), predating #1657 and predating the branch that found it. NOT fixed in flow: it changes the refusal ordering of both arms and wants its own red-first test per arm. Owned by [`dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) `## Owed` O31 | bug | -| [#1131](https://github.com/mudler/vllm.cpp/issues/1131) | `MUSIC3-DIT-ARM-REACH` | MiniMax-Music3's 2.4B fp32 DiT reaches the device through a switch **no gate could enter**, so the capability could be deleted and every suite stayed green. Two mutations on the shipped code measured it: `on_device = false` in `Music3DenoiseChunks` -- the production call site disabled -- left `test_minimax_music3_acoustic` 32/283, `_speech` 9/223 and `test_speech_engine` 11/38 GREEN and unchanged, and `Music3DenoiseDeviceArm::half_set() -> false` left the same three GREEN. `DitForwardDevice` and `StageMusic3DitWeights` were gated as a CLASS the whole time, against upstream's own goldens on both arms; what nothing reached was the WIRING that turns them on, so a change that silently returned the DiT to the host loops would have produced a correct song hours late and no red. STRUCTURAL, not careless: the engine selects on `queue_.device.type != vt::DeviceType::kCPU`, and `src/vllm/multimodal/speech_engine.cpp::SpeechEngineDeviceType` refuses `--speech-device 1` on a CPU-only build before a queue exists, so on every runner CI owns that condition is false and a branch written at that line is unreachable by construction. Fixed in flow by this row, the same way `minimax-music3.md` §19.5 fixed the depth arm's twin: the RULE moves out of the engine into `Music3SelectDitArm`, which runs on both sides of the condition and is therefore drivable from a CPU gate with a real CPU queue and with a fabricated non-CPU one that must engage or refuse BY NAME -- the silent host fallback being the third outcome and the defect -- and the SELECTED arm is then driven through `Music3DenoiseChunks`, the production function the engine calls, with the gate asserting WHICH ARM RAN from an exact call count on the production `denoise.dit_device` profile bucket and on `dit.pack` inside `DitForwardDevice`, never from output agreement, because the two arms agree by design. Both of the mutations above now red. NOT fully closed and the issue stays OPEN: the engine's own one-line CALL to `Music3SelectDitArm` still cannot be gated on a CPU runner -- it needs the 28.5 GB checkpoint and a real accelerator -- which is the identical residual `minimax-music3.md` §19.7 carries for the depth arm, is recorded under `## Owed` in `.agents/specs/music3-dit-arm-reachability.md`, and closes with a `thor:gpu0` leg under an `rc` lease | bug | -| [#1674](https://github.com/mudler/vllm.cpp/issues/1674) | `ENG-PUBLIC-DOC-PROJECTIONS` | Retire the shared STATUS projection, split public benchmark records, and remove documentation policy gates from lifecycle changes | policy | -| [#1679](https://github.com/mudler/vllm.cpp/issues/1679) | `GATE-QWEN38-27B-FP8-BLOCK` | **The pinned oracle selected `FLASH_ATTN` on GB10 and then loaded and generated on it, against a MEASURED record that says this wheel cannot.** Measured on `dgx:gpu0` (NVIDIA GB10, driver 580.173.02, compute capability 12.1, aarch64) on 2026-08-22, `rc` job `0d5dfa6a-195f-4475-8527-538ad91102c8`, logs on the share at `/mnt/nas_share/rc/fp8-oracle/`. Identity was asserted before anything else, from `/` so no source tree could be imported by accident: `vllm.__file__ = /tmp/oracle-venv/lib/python3.12/site-packages/vllm/__init__.py`, `vllm.__version__ = 0.1.dev1+g555967922` carrying the parity pin `5559679229bc961848b121ccdeaa8fa5d79bec98`, `torch 2.13.0+cu130`, `IDENTITY_RC=0`. The run exported `VLLM_ATTENTION_BACKEND=FLASHINFER` and the engine logged `Using FLASH_ATTN attention backend out of potential backends: ['FLASH_ATTN', 'FLASHINFER', 'TRITON_ATTN', 'FLEX_ATTENTION']`, then loaded `/workspace/ckpt/qwen38-27b-fp8` and generated 24 greedy tokens with no attention-backend failure. `/mnt/nas_share/rc/oracle-vllm/README-WHEELS.md` records as a MEASURED claim that this wheel's FlashAttention carries no sm_12x code, that `vllm-project/flash-attention` hard-codes `FA2_ARCHS "8.0+PTX"`, and that reaching sm_121 therefore needs a driver JIT of compute_80 PTX which fails with `cudaErrorUnsupportedPtxVersion`; [#1456](https://github.com/mudler/vllm.cpp/issues/1456) reaches the same conclusion from a from-source build at another revision, reporting `_vllm_fa2_C.abi3.so: sm_80` and `_vllm_fa3_C.abi3.so: sm_75`. It did not fail here. **NEITHER SIDE IS ASSERTED.** The README is not declared wrong from a log line, and `FLASH_ATTN` is not declared safe on GB10 because one decode completed; the reason is unknown and the two records disagree. Candidates not resolved here, each a different fix: the selector reached a path that never entered an FA2 kernel; the FA2 entry point was reached and the driver JIT succeeded where the README measured it failing; or the installed wheel differs from what the README describes. One part of the run is already known to be wrong and explains nothing about the above -- the same README records that `VLLM_ATTENTION_BACKEND` and `VLLM_FLASH_ATTN_VERSION` do NOT exist at this pin and that the backend is chosen by the `LLM(..., attention_backend=...)` argument, so the probe may never have asked for FLASHINFER through a lever the pin honours, which is a recipe defect the gate run must not repeat. It matters because a gate pins the executed backend on BOTH sides, and the oracle's selection on this box is not doing what the environment asked. What would settle it: one lease that constructs the engine with the `attention_backend` argument rather than the variable, records which backend each construction selects, and reads the installed wheel's FA modules with `cuobjdump --list-elf` so the sm_12x claim is checked against the ARTIFACT; do not close it from a log alone. NOT fixed in flow: the run that found it was a feasibility probe of the ORACLE side only, it produced no gate result, and adjudicating an attention-backend selector needs its own lease and its own measurement. Found while measuring the oracle-refusal stop condition of [`gate-qwen38-27b-fp8-block.md`](../specs/gate-qwen38-27b-fp8-block.md), which lists it under `## Owed` | verification | -| [#1664](https://github.com/mudler/vllm.cpp/issues/1664) | `VT-CONV1D-TIME-BLOCK` | **The MiniMax-Music3 vocoder's parallel decomposition returns 4.27x of 14 threads and nothing has measured why.** `.agents/specs/minimax-music3.md` §18.8b measured the decode window at **2.16x per core** on `thor:gpu0` and **1.365x at 14 threads** -- a scaling of 6.76x before the tiling and 4.27x after, of a possible 14x, so the FASTER arm scales WORSE -- and attributed it to a shared-bandwidth limit **while naming that as an inference**: *"no bandwidth counter was read, and none is available on this worker"*. `vocoder.decode_window` is **122.169 s of a 449.969 s run** at 20 s / 30 steps, about 27 %, and it is the largest term no row owns. `Conv1dKernel` (`src/vt/cpu/cpu_conv1d_general.cpp:152`) partitions `rows = batch * out_channels` and nothing else, so every thread sweeps the WHOLE input tensor for its own slice of output channels and the reuse available across output channels is never taken: at 344 latent frames the b3 residual activation is 67.6 MiB against 96 output channels, and the b1 residual 33.8 MiB against 384. Five candidates are separable by ablation and none had been separated -- activation residency, `Threadpool::Barrier` cost over too many small parallel regions, work-partition granularity (`conv_out` has ONE output row, so it runs INLINE on the caller at every thread count), the pool not running, and the CPU clock falling as cores light up. `ncu` is refused on this fleet (`ERR_NVGPUCTRPERM` even from the root `rc` worker) and Thor is aarch64, so the instrument is ablation plus `/sys` clock sampling and an unavailable counter is reported as unavailable rather than replaced by an estimate. Any repair is bound by BIT-IDENTITY: `vt::Conv1d` is bit-identical to its serial reference by construction and four models decode through it -- MiniMax-Music3, LTX-2.5's audio VAE (13 arms at `5e-6`, the tightest), IndexTTS-2.5's BigVGAN, MiniMax-H3's audio VAE -- so a decomposition that moves which thread computes which cell is admissible and one that moves a per-cell summation order is a different row. Spec [`vt-conv1d-time-block.md`](../specs/vt-conv1d-time-block.md); parent lane [#672](https://github.com/mudler/vllm.cpp/issues/672), owed under §18.9 of [`minimax-music3.md`](../specs/minimax-music3.md) | perf | -| [#1683](https://github.com/mudler/vllm.cpp/issues/1683) | `VT-CONV1D-TIME-BLOCK` | **The 11.48x scaling curve was measured on arm C, which does not ship, and the 4.11x ratio is composed across two jobs, one of which the row declares defective.** Found by the fresh review of [PR #1678](https://github.com/mudler/vllm.cpp/pull/1678) at head `e1396fc73`. `.agents/specs/vt-conv1d-time-block.md` §2b's scaling table carries the column header `arm C` and §2b defines arm C as `cf9296496`, "B + the conv decomposition, blocked UNCONDITIONALLY"; the shipped tree is arm D (`0f738d6ec`), which added the `out_channels * kernel <= in_len` condition in `06ba79d1b` and was measured at exactly ONE operating point -- 86 latents, 14 threads (§2c). The 1/2/4/8/14 sweep that produces 11.48x was never run on the shipped arm, and the number was nevertheless attributed to it in four places: the PR title, spec §9, `docs/STATUS.md` and `docs/BENCHMARKS.md`. `.agents/benchmark-record.md` got it right -- it prints the `arm C` header and names C as unconditional -- so the forensic record knew and the public projections dropped the label. The error is very likely CONSERVATIVE (at a 20-latent window the condition declines the two b0 shapes where C measured 0.82x/0.89x, so D >= C) but that is an inference, and the row's own standard is that an unmeasured quantity is reported as unmeasured. SEPARATELY, arm D's median 3.4989 s comes from job `214f5f70` (300 s settle, seven alternated rounds) while the 14.3895 s denominator comes from job `3ca07477`, of which the spec itself says the rounds "ran at `uptime` load 8.84, the decaying residue of three back-to-back builds in the same lease" and records that as "a defect in the SCHEDULE of the job rather than as a result" -- same boot id and worker, but a different job, not alternated against D and not under the settle §2c exists to provide, and whether arm A's 86-latent leg fell inside that window is not stated and cannot be recovered from the tree because no job log is committed. The 1.067x B-against-D figure beside it IS properly paired and is unaffected. FIXED IN FLOW as far as labelling goes -- all four public sites now name the curve as arm C's, say the shipped arm has one measured point, and name the denominator's contention. NOT fixed as a measurement: it needs one `thor:gpu0` lease, one boot id, both arms built inside it with distinct binary sha256, a settle with `uptime` on both sides, and 1/2/4/8/14 swept with the arms ALTERNATED at each count, which produces the shipped arm's own curve and a paired A-against-D ratio in one job. Owned by [`vt-conv1d-time-block.md`](../specs/vt-conv1d-time-block.md) `## Owed` | verification | -| [#1684](https://github.com/mudler/vllm.cpp/issues/1684) | `VT-CONV1D-TIME-BLOCK` | **No per-MODEL suite exercises a `blocks > 1` shape, so a defect confined to the time-blocked axis reds the op's suite and nothing else.** Measured by the fresh review of [PR #1678](https://github.com/mudler/vllm.cpp/pull/1678) rather than argued: mutation M7b sign-flips every output cell of `Conv1dKernel` when `blocks > 1` -- exactly and only the axis the row adds -- and at head `e1396fc73` EIGHT of the ten suites stayed GREEN (`test_vocoder1d`, `test_bigvgan`, `test_minimax_music3_acoustic`, `test_ltx2_vae`, `test_minimax_h3`, `test_indextts2_pipeline`, `test_indextts2_family`, `test_ops_conv1d_depthwise`), with only `test_ops_conv1d_general` and `test_host_parallel` red. All four audio consumers reach `vt::Conv1d`; every one of them does so at SINGLE-BLOCK shapes only. The row disclosed the gap in §6c and compensated with a geometry gate that asserted `blocks > 1` on the vocoder's shapes -- but that gate read SIX hard-coded shape constants transcribed by hand from `minimax_music3_loader.h:253-265`, so a loader change would have left it green while the production shapes collapsed to one block. PARTLY FIXED IN FLOW, both halves: the geometry gate now DERIVES its shapes by walking `MiniMaxMusic3VocoderConfig` and `kVocoderResidualDilations`/`kVocoderResidualUnits`, and the derivation is proved by mutation -- setting `upsampling_ratios{1,1,1,1}` reds the derived gate (1 case / 16 assertions, rc 1) and left the hand-transcribed one at 14 cases / 19 615 assertions / `SUCCESS!`, rc 0, completely blind; and `tests/vllm/models/test_vocoder1d.cpp` gained `vocoder1d Conv1d is exact ACROSS a time block boundary`, which enters through `vllm::vocoder1d::Conv1d` -- the body all four audio models call -- at 32 channels, kernel 7, 10 000 positions, asserts the block length is shorter than the output length so the case cannot silently become single-block, and reds under M7b at the repaired head (1 case / 4 assertions, rc 1). STILL OPEN: the four MODEL suites reach the provider at single-block shapes only, so M7b still leaves them green. Closing it means lengthening each consumer's reduced-dimension fixture until its convolutions cross a block boundary, which moves those fixtures' goldens -- a fixture change per model, not a test addition. Owned by [`vt-conv1d-time-block.md`](../specs/vt-conv1d-time-block.md) `## Owed` | verification | -| [#1685](https://github.com/mudler/vllm.cpp/issues/1685) | `SPEC-DFLASH2` | **The DFlash2 speed run's denominator declares `TRITON_ATTN` while five of its layers resolve `FlashAttentionBackend`, which `FA-CONSTRAINT.txt` says cannot target sm_12x.** Observed by the run that produced the first DFlash2 speed ratio (`dgx:gpu0`, `rc` job `ec9cf6cd-0aaf-4323-806d-6a12da2bd08f`, 2026-08-22, gate tree `d25730fbb`, `GATE_RC=0`, ours/vLLM = 0.8016987337853048 on `output_throughput_tok_s`). The gate did NOT refuse; this is a question about what the denominator IS, not a gate failure. `evidence/vllm-arm.json` records the scalar `attention_backend: TRITON_ATTN` from `attention_backend_source: read_back_from_engine`, and beside it the per-group census read off `...model_runner.attn_groups`: `GDNAttentionBackend` 48 layers, `TritonAttentionBackend` 16, and `FlashAttentionBackend` 5 -- the five being `model.layers.64-68.self_attn.attn`, i.e. the DFlash2 draft's sliding-window layers. `/workspace/oracle-dflash2/FA-CONSTRAINT.txt` records `FA_USABLE=0` for sm_12x from [#1456](https://github.com/mudler/vllm.cpp/issues/1456), where a real source build at `CUDA_ARCHS=12.0` emitted `sm_80` for `_vllm_fa2_C` and `sm_75` for `_vllm_fa3_C`, and `TRITON_ATTN` is the declared backend on this box BECAUSE of that. The engine nevertheless loaded, generated 64 tokens on every one of its 20 legs, and produced a clean 83-sample clock window. Three readings and none established: upstream intends the draft's sliding-window layers to select FA independently of the engine-wide choice; the label over-describes what executes and those layers fall back at runtime; or FA genuinely runs on sm_121 through driver JIT of the emitted `sm_80` PTX, in which case #1456's CONCLUSION -- not its measurement -- needs re-reading. This is the THIRD independent observation of FA resolving on this box against `FA_USABLE=0`. Settled by reading the pinned wheel's own selection path for sliding-window and draft layers and tracing whether those five layers dispatch FA kernels or fall back; either outcome is cheap and changes what the scalar should say. Recorded under `## Owed` O26 residual 1 of `.agents/specs/dflash2-spec-decode.md` and in the 2026-08-22 entry of `.agents/benchmark-record.md`. Related: [#1658](https://github.com/mudler/vllm.cpp/issues/1658) (why the per-group map is recorded beside the scalar at all: one string cannot describe 48 + 16 + 5) and [#1673](https://github.com/mudler/vllm.cpp/issues/1673) (the other open caveat on the same ratio). Evidence: `/mnt/nas_share/rc/dflash2-1673/out-n1673b/evidence/{vllm-arm.json,clock-vllm.json}` and `out-n1673b/m-gate.log`. | bug | -| [#1681](https://github.com/mudler/vllm.cpp/issues/1681) | `SERVE-CHAT-TEMPLATE` | `POST /v1/chat/completions` answers HTTP 500 for the whole Qwen3.8 family because the vendored minja Jinja engine implements twelve of Jinja2's built-in tests and `undefined` is not one of them, so `{%- if enable_thinking is undefined or enable_thinking is true %}` throws at row 46 of the checkpoint's own template. `is true` was already present, so the first term was the only break. Fixed in flow together with the second half of the same defect: `enable_thinking` was set unconditionally by `apply_chat_template`, so even with `undefined` implemented the variable could never be undefined and the Qwen3.8 default would have been thinking-OFF against upstream's thinking-ON, and `ChatCompletionRequest` carried no `chat_template_kwargs` at all, so the `{"chat_template_kwargs":{"enable_thinking":false}}` body both competitor arms of [#1574](https://github.com/mudler/vllm.cpp/issues/1574) were measured with was silently ignored. Spec [`chat-template-jinja-undefined.md`](../specs/chat-template-jinja-undefined.md), whose `## Owed` carries the twelve of Jinja2's thirty canonical built-in tests that stay unimplemented, each with the reason: nine need a grammar change because minja parses the right side of `is` as a bare identifier, `filter` and `test` need a name registry minja does not have, and `callable` can never be handed a callable because `BinaryOpExpr::do_evaluate` defers every binary operation whose left operand is one. No chat template of any checkpoint in `docs/USAGE.md` uses any of the twelve | bug | -| [#1625](https://github.com/mudler/vllm.cpp/issues/1625) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **Captured Tenstorrent decode hangs deterministically on the first MULTI-request run, while every single-request captured leg and the whole host-free eager path work.** Reproduced twice on the P150 at tree `b86e3705f` (main `52e328789` + the R5 flip): `VT_DUMP_IDS=1 test_qwen3_paged_engine` (16 sequential requests) stalls ~10 s into stepping with one tt-metal worker spinning at 100% and the main thread blocked, killed after 11 min; the last device log line is the allocator warning `Allocating device buffers is unsafe due to the existence of an active trace` (allocator.cpp:123). `VT_TT_RECAPTURE_EVERY=8` (live traces destroyed every 8 replays) hangs IDENTICALLY at the same point, ruling out a per-trace replay-count cap, and `VLLM_CPP_CUDAGRAPH=0` (host-free eager, no capture) completes the same gate in 35 s at 125/125 assertions and 10.94/10.95/11.06 tok/s warm. NOT FIXED IN FLOW: the mechanism is undiagnosed (eager-alloc-around-live-trace across the request boundary is the recorded hypothesis class, qwen3.cpp Step() comment, but the RECAPTURE_EVERY result narrows it), so #1604 lands with `support_static_graph_mode()` declined by default on TT (opt-in `VT_TT_DECODE_CAPTURE`); flipping capture back on by default is owned by this issue | bug | -| [#1626](https://github.com/mudler/vllm.cpp/issues/1626) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`test_mistral_paged_engine.cpp:102`'s `RunGate` takes `const char* label` and streams it into doctest MESSAGE/REQUIRE_MESSAGE, so every label renders as `1`.** The #1604 R5 dump run printed `1 dumped our token ids -> ...` and the anchor-drift REQUIRE printed `logged: 1 anchor drift prompt[3] ...`. Same defect class #1508 fixed in the Qwen3 gate (doctest `MessageBuilder` has no `const char*` overload); the Mistral copy never got the `const std::string&` fix. FIXED IN FLOW by the same #1604 change (one-line parameter change, call site converts implicitly) | bug | -| [#1627](https://github.com/mudler/vllm.cpp/issues/1627) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **The Tenstorrent backend has no `SupportsAsyncSampledTokenReadback` override, so async scheduling resolves OFF on TT and `test_qwen3_dense_async_serving` FATALs on every cached checkpoint** (3 FATAL / 5 checkpoint-absent skip on the P150) at the anti-vacuous-pass guard `REQUIRE(loaded->async_scheduling_enabled())` (`test_qwen3_dense_async_serving.cpp:124`). Mechanism: `runner_supports_async()` derives from `vt::Backend::SupportsAsyncSampledTokenReadback()` (`runner.cpp:109-112`, default false at `backend.h:186`), overridden only by CPU and CUDA. PRE-EXISTING: zero hits under `src/vt/tenstorrent/` at base `52e328789`, and the R5 flip commits touch none of the resolution path — captured-vs-eager decode mode is orthogonal. NOT FIXED IN FLOW: enabling it needs a device-mirrored sampled-id design against the tt-metal allocator (CUDA's `async_device_mirror` equivalent) plus the #323-class guard re-proven on device — its own spec and gates; owned by this issue | bug | -| [#1688](https://github.com/mudler/vllm.cpp/issues/1688) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`ReshapeAndCacheKernel` latched `VT_TT_HOST_FREE_DECODE` in a function-local `static`, so after the R5 default flip the documented opt-out `VT_TT_HOST_FREE_DECODE=0` silently did not reach that path.** Found by the fresh review of [PR #1630](https://github.com/mudler/vllm.cpp/pull/1630) at head `450a524b0`. `tenstorrent_device.h` documents `HostFreeDecodeEnabled()` as "No function-local static caching: tests toggle this env per case in one process", and the flip converted eleven call sites to a live read — `EnsureDevice2D`, `RmsNormKernel`, `PreferDeviceRope`, `CopyDeviceDeviceIfCapture`, `MemsetDeviceIfCapture`, `WarmRopeCosSin`, `WarmPagedKvShadow`, `WarmRacIdx`, `WarmPaMeta`, `WarmDecodePos`, `CaptureDecodePosAdvance` — but left `tenstorrent_ops.cpp:2130` a `static const bool`. The polarity flip is what made it bite: pre-flip the latch cached the default-OFF state so only the opt-IN could be defeated, and that was set before the process started; post-flip it caches default-ON, so the opt-out that the flip tells operators to use never arrives at RAC once any decode has run. The suite could not see it — the flip's own `support_static_graph_mode` round trip and the `setenv(...,"0")` inertness guard both prove the contract elsewhere and neither reaches `ReshapeAndCacheKernel`. FIXED IN FLOW: the `static` is dropped, matching every other converted site. NOT fixed as a test: `ReshapeAndCacheKernel` needs a real Blackhole device, so every case reaching it is behind `TenstorrentPresent()` and skips on every `rc` fleet host; the owed `thalia` case is recorded under [`tenstorrent-host-free-forward.md`](../specs/tenstorrent-host-free-forward.md) `## Owed` | bug | -| [#1661](https://github.com/mudler/vllm.cpp/issues/1661) | `TEST-SHELLCHECK-SKIP-GUARD` | **`test_script_stays_shellcheck_clean` ERRORs instead of skipping on hosts without the shellcheck binary: the guard probes `returncode`, but a missing binary raises `FileNotFoundError` before a returncode exists**, so the `tools suites` preflight gate is red on such hosts (`thalia`, aarch64). The test predates the window (last touched `b95543c44`); it surfaced when current main's preflight started running the module. FIXED IN FLOW by the same change: the probe gains the `except FileNotFoundError` skip arm (message unchanged), and both arms are proven — no-binary host skips via the new arm, a PATH-shim `shellcheck` exiting 3 still skips via the returncode arm, and unmodified main reds on this host | bug | -| [#1709](https://github.com/mudler/vllm.cpp/issues/1709) | `LTX25-DIT-ATTN-FLASH` | **`dgx:gpu0` holds 110.41 GiB that belongs to no `/proc/meminfo` category, and `rc` keeps handing out leases against it.** Measured read-only from inside a lease (`rc` job `ab12aac1-b862-4ac6-8292-9f2c641e6a8d`, 2026-08-22T16:55:46Z), because the fleet rule forbids `ssh`: `MemTotal` 119.63 GiB, `MemFree` 5.23, `AnonPages` 0.93, `Cached` 0.91, `Buffers` 0.05, `Shmem` 0.04, `Slab` 1.10, `VmallocUsed` 1.00 -- **accounted 9.22 GiB, UNACCOUNTED 110.41 GiB, 92.3% of the box.** Every ordinary owner is excluded by measurement and not by argument: the sum of `VmRSS` over every visible `/proc/*/status` is 0.0 GiB; `/dev/shm` is a 64 M mount containing nothing; `nvidia-smi` reads 0%, 11 W, `No running processes found`; the cgroup reports `memory.max=max` and `memory.current` 113 MiB; and the value is 5.0 GiB at 15:49Z, 5.1 GiB flat through 16:39Z and 4.98 GiB at 16:55Z across four leases by three submitters, on a box `up 2:33` at load 0.25. The leading explanation is a driver-held unified-memory allocation that outlived its process, which fits every observation including `Memory-Usage: Not Supported` being the one meter that would have named it -- recorded as a HYPOTHESIS, since confirming it needs host access. COST: three leases on #1612 -- `5fb9399f` lost its worker to an OOM during a build started against 5 GiB, `2ccd1acf` waited its full 1200 s at a flat 5.0 GiB and refused with exit 39, `ab12aac1` measured the box. FIXED IN FLOW, the half this row owns: `scripts/ltx25-dit-attn-flash-pixel-ab.sh` gains a `MemAvailable` start gate that waits, logs EVERY poll and refuses by name -- and logging every poll is the only reason the condition can be called persistent rather than busy. NOT FIXED IN FLOW, and named rather than folded in: a device-readiness condition in the controller, which would have parked every one of these jobs instead of spending them, and a line in the DGX profile saying a granted lease does not imply a reclaimed box. Owner: row `LTX25-DIT-ATTN-FLASH`, under `## Owed` in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) | bug | -| [#1689](https://github.com/mudler/vllm.cpp/issues/1689) | `DOCS-MODELS-HUMAN` | **All 13 pages under `docs/models/` are written for a reviewer of the change that produced them, not for a reader who wants to run the model.** Three shapes recur. Twelve of the thirteen open with the identical line `Use this page for checkpoints, commands, supported arms, and current limitations`, which says nothing the title did not. Defect archaeology precedes the command: `qwen3-8-2-4t.md` was 593 lines and spent its first 100 on [#1123](https://github.com/mudler/vllm.cpp/issues/1123), [#1124](https://github.com/mudler/vllm.cpp/issues/1124) and [#1299](https://github.com/mudler/vllm.cpp/issues/1299) before saying what to type, reaching the serve command at line 434. And spec content sits in a public projection: rationale for a default and the history of a fixed defect belong in the row spec. FIXED IN FLOW: every page now leads with what the model is, what it costs, the exact command, and the honest limits, in that order. **No measured value changed** — a string-level retention check asserts that every checkpoint revision, sha256, byte count, decode figure, refusal message and issue link on the old `qwen3-8-2-4t.md` still resolves in the tree. The mechanism content removed from that page was NOT deleted: `docs/guides/expert-streaming.md` was a 17-line stub that the model page already claimed owned the config schema, precedence rule, statistics line and per-device conditions, so that material moved there and the stub became the owner it was described as. `docs/models/README.md` becomes a real index naming what each page answers. Related: [#1691](https://github.com/mudler/vllm.cpp/issues/1691), the stale container claim found by the same work | record | -| [#1691](https://github.com/mudler/vllm.cpp/issues/1691) | `DOCS-MODELS-HUMAN` | **`docs/QUICKSTART.md` said no container lane had ever been published, and three had.** The page carried `The container lanes have never been published, so no tag below resolves against the registry yet` and, on its `docker run` line, `The package carries only a stage tag today, which is a build artifact and not a lane`. Both were true when written and false on 2026-08-22: `ghcr.io/mudler/vllm.cpp` is public and carries `main-cpu`, `main-cuda` and `main-vulkan`, each a `linux/amd64` + `linux/arm64` manifest, published 07:30 UTC that day. `docs/guides/container-images.md` already documented the `:main-` tags, so the two pages disagreed. VERIFIED by execution rather than by reading the registry listing: `docker run --rm ghcr.io/mudler/vllm.cpp:main-cpu --version` answered `vllm.cpp 0.0.3 c-abi=23` at digest `sha256:7f88301ea282dad778748929e7aa6869d2418c8d295eef0e7900cca8310d06e5`, and the same image with a mounted `Qwen/Qwen3-0.6B` returned tokens through `/v1/completions` on host `mudler-ubuntu-box` (x86_64, Docker 29.1.2). The image also parses the `vllm_cpp` weight-residency document and echoes `mmap=on prefault=off expert_stream=on expert_stream_slots=4000`, which is what lets the Qwen3.8 2.4T page carry a container form of its recipe. FIXED IN FLOW: the false note is corrected, the `docker run` line names a tag that resolves, and the executed-row table gains its first real row. **This does NOT close [#1281](https://github.com/mudler/vllm.cpp/issues/1281)**: `:latest` still does not exist, `--model org/repo` is still blocked by [#1511](https://github.com/mudler/vllm.cpp/issues/1511) so the executed row mounts a local directory, and no GPU-lane row was run. The executed-row obligation stays owned by #1281 | bug | -| [#1690](https://github.com/mudler/vllm.cpp/issues/1690) | `DOC-README-ABI-LANDING-SOURCE` | **`check-doc-checkpoint.py`'s `LANDING_SOURCE_FILES` omitted `include/vllm.h`, so a commit that bumps `VLLM_ABI_VERSION` could never repair the README claim it invalidated.** Found by the review of [PR #1655](https://github.com/mudler/vllm.cpp/pull/1655). The README `## Use it as a library (C API)` block quotes the ABI version out of the header; the header was in `USER_USAGE_FILES` (so an ABI change owes `docs/USAGE.md`) but not in `LANDING_SOURCE_FILES`, and the README rule refuses a claim change unaccompanied by a landing source. The claim could therefore be invalidated but not repaired by the same edit, which is how the README reached `VLLM_ABI_VERSION 21` against a header reading `23`, alongside a stale "46 exported functions" for a header declaring 47. That second half is stale by one rather than by six -- an earlier review figure of 51/52 swept in the `#define VLLM_API` visibility block and counted `vllm_*` identifiers that are typedefs and struct fields rather than exported functions -- so the case for deleting the count is that a live count of one file stored in another goes stale on any ABI addition, not that it is badly wrong. The set's own criterion already admitted it -- the checker's comment says every member is "something the README QUOTES" -- and the header was the only such source missing. FIXED IN FLOW: `include/vllm.h` is added to the set, red-before/green-after pinned by `test_the_c_abi_header_is_a_landing_source`, with `test_the_c_abi_header_permits_but_does_not_demand_readme` proving no new README obligation and the pre-existing no-class tests still green | bug | -| [#1702](https://github.com/mudler/vllm.cpp/issues/1702) | `LTX25-DIT-ATTN-FA2-HD128` | **Three of the four LTX-2.5 artefacts every render is fed have no row in the `docs/USAGE.md` checkpoint registry.** AGENTS.md requires every ported model to document the checkpoints it was built and gated against there -- file name, size, exact HuggingFace repo AND revision, sha256 for a quantized artifact -- in the same change that makes the capability reachable, because weights are the one part of a port a reader cannot infer from the code. The registry carried four LTX-2.5 rows and all four are DiT or LoRA, while `ltx2-gen` takes four artefacts on four flags: `--video-vae` (`ltx-2.5-video-vae-conv-bf16.safetensors`), `--audio-vae` (`ltx-2.5-audio-vae-bf16.safetensors`) and `--encoder` (`gemma4-12b-with-proj-nvfp4-torchao.safetensors`) had NO row, and the encoder is NVFP4 torchao -- the quantized case the rule singles out for a sha256, and a file the loader already refuses by name when malformed (`ltx2_text_encoder.cpp`). Every LTX-2.5 job in the record feeds exactly these, including the runs behind the 47.84 s and 7.680 s numbers, so three files nobody can identify from the repository sit under every LTX measurement it has. FIXED IN FLOW by `LTX25-DIT-ATTN-FA2-HD128` ([#1551](https://github.com/mudler/vllm.cpp/issues/1551)), which is the change that makes the DiT reachable on the FA-2 tensor-core path and is gated against these exact four files. Provenance READ, not inferred: `huggingface_hub` writes a `.metadata` sidecar per download holding `commit_hash`, `etag`, `timestamp`, and all four `Lightricks/LTX-2.5` sidecars carry the SAME `commit_hash`, which is what makes it a snapshot revision rather than a blob id. The etag of an LFS file is its sha256 and all three were CONFIRMED by hashing the local bytes rather than trusted, because an unauthenticated HuggingFace tree API has returned a fabricated content hash for a gated repo here before; all three matched. NOT fixed, and named rather than folded in: the distilled NVFP4 DiT row pins `6c7e5e57...` while that file's own sidecar records `8a4ff96f...` -- both can be true, and choosing without knowing which bytes that row's author measured swaps a possibly-stale pin for a definitely-unverified one. That half is under `## Owed` in [`ltx25-dit-attn-fa2-hd128.md`](../specs/ltx25-dit-attn-fa2-hd128.md) | bug | -| [#1692](https://github.com/mudler/vllm.cpp/issues/1692) | `KERNEL-ACCEL-PROVIDER-DECLINE-EXACT` | **[#1584](https://github.com/mudler/vllm.cpp/issues/1584) is FIXED by this row, and its gate is CPU-only: both production call sites need a GPU and neither arm was executed.** The fix adds `GetOpFallbackUncounted`, sharing one body with `GetOpFallback` so resolution order, the reference-tier install, the drain and every throw stay identical and only the `fetch_add` is conditional -- in its existing position, BEFORE the "nothing below" check, because a decline that throws is still a decline. The two SHAPE-GATED providers that hoist the lookup into a function-local static switch to it: `cuda_attention_cross.cu` `BlockedFallback` and `metal_mlx_provider.mm` `MlxFallback`. **Candidate 1 of #1584 -- drop the count from `GetOpFallback` -- was REJECTED and measured rather than argued:** it edits five per-call sites (`vulkan_ops.cpp:950,1067,1488,1509`, `tenstorrent_ops.cpp:1341`), one already gated on an exact value at `test_vulkan_backend.cpp:2901` (`after.declines == before.declines + 1`), and its failure mode for a future caller is silent UNDER-counting, the Risk 4 the seam exists to expose. Net blast radius is ZERO backends, not the four #1584 estimated; exactly two providers change what they report, by one. #1555's `WarmDeclineOnce` workaround is REMOVED. **What this issue owns is the missing EXECUTION:** `test_ops_attention_cross` on CUDA (20 cases, ALL 20 SKIP on a CPU-only build -- 32 assertions, every one the skip guard, so that suite gives the change no coverage here), `test_metal_backend` on a `VLLM_CPP_MLX` build, and the `.agents/reachability.md` mutation on `BlockedFallback()` / `MlxFallback()`. The CUDA COMPILE is covered by the `-DVLLM_CPP_CUDA=ON` `vllm`-target job in `.github/workflows/ci.yml`; the Metal `.mm` file is compiled by NO job in this repository, because MLX needs `MLX_ROOT`. Red-before/green-after and the mutation table in [op-provider-decline-exact.md](../specs/op-provider-decline-exact.md) | bug | -| [#1718](https://github.com/mudler/vllm.cpp/issues/1718) | `LTX25-PHASE-INSTRUMENT` | **`PhaseLog`'s instrument-charge attribution is gated only in AGGREGATE: deleting any ONE of the five charge sites leaves the whole suite green.** The sites in `src/vllm/multimodal/render_phase_log.cpp` are `Open`'s pre-lock mutex wait (the `entered` clock read, which the row's `## Design` 1 names as the reason the mechanism exists), `Open`'s tail after `SampleLocked`, `Close`'s tail, the sampler join, and `SampleLocked`'s own self-charge. Measured by mutation at the head of [PR #1711](https://github.com/mudler/vllm.cpp/pull/1711), each printing `compile_status`, `git diff --numstat` and the doctest `test cases:` / `assertions:` lines, tree restored from a pristine byte copy and verified by sha256: `N6` (`Open`'s pre-lock charge replaced by a `(void)` cast of the same expression) GREEN, `N4` (`ChargeLocked`'s `if (from < 0.0) return;` becomes a clamp to zero, which the comment beside it argues at length is the defect that makes a gate pass) GREEN, `NNOSORT` (`ByStart` stable-sorts an empty range, i.e. does not sort at all) GREEN. Only `R2`, which makes `ChargeLocked` return immediately so nothing anywhere is charged, reds -- at `REQUIRE(instrument > 0.0)` in the render case. `NREVSORT` beside `NNOSORT` says exactly how much of `ByStart` is held: a DESCENDING sort reds at `CHECK(seconds >= 0.0)` on -16.996 ms and a REMOVAL does not, because every timeline this suite builds is already start-ordered. N6's first staging is itself worth recording: written as `if (false) { ... }` it left `entered` unused, failed `-Werror`, and printed `compile_status=1` -- a mutation that fails to build reads exactly like a passing test and only the printed status separated them. NOT FIXED IN FLOW, and the reason is measured rather than asserted: a per-site case was written and run THREE times and does not measure the site. `PhaseLog::Records()` is the only public entry point that holds the process-wide mutex without charging itself, which is what makes it usable as a hold at all, and contention staged through it lands mostly in `PhaseLog::Close`'s lock wait -- charged to nobody ([#1719](https://github.com/mudler/vllm.cpp/issues/1719)) and inside the CHILD's duration rather than the parent's charge. The case passed in ISOLATION at a separation of 615x and failed 5 of 5 inside the suite, twice because no reader held the lock at all and three times because a contended parent that took 21 ms was charged 112 us. A flaky gate over an instrument whose whole subject is flaky gates is the failure this cluster exists to stop, and [#1569](https://github.com/mudler/vllm.cpp/issues/1569) is the standing record of what an instrument whose own mutation cannot fail costs. What would settle it: a probe that holds the mutex through an entry point that does not charge itself AND lands the wait on the site under test, which needs #1719 first. Owned by [`ltx25-phase-instrument.md`](../specs/ltx25-phase-instrument.md) `## Owed` | verification | -| [#1719](https://github.com/mudler/vllm.cpp/issues/1719) | `LTX25-PHASE-INSTRUMENT` | **`PhaseLog::Close`'s pre-lock wait is charged to nobody, so it inflates the closing record's UNCOVERED time -- the quantity the coverage floor reads.** `PhaseLog::Open` reads a clock BEFORE it takes the process-wide mutex (`const std::chrono::steady_clock::time_point entered = std::chrono::steady_clock::now();`, `src/vllm/multimodal/render_phase_log.cpp`) and charges that lock wait to whatever encloses the gap before the new record. `Close` has no such read: it takes the mutex and only then stamps `r.end`, so the wall between a scope's last statement and its recorded end is instrument time that no record and no table absorbs -- not in `Record::instrument_seconds`, not in `PhaseLog::Instrument()`, and therefore not in `unaccounted_seconds`'s explanation. It is not cosmetic: that interval lies inside the CLOSING record's own duration and outside every child of it, which is exactly what `covered >= min_coverage * leaf_seconds` reads -- the floor [#1439](https://github.com/mudler/vllm.cpp/issues/1439), [#1470](https://github.com/mudler/vllm.cpp/issues/1470), [#1494](https://github.com/mudler/vllm.cpp/issues/1494) and [#1536](https://github.com/mudler/vllm.cpp/issues/1536) have spent months arguing about. A `Close` that waits on a contended mutex therefore makes that floor see an un-named phase, and the number a reader would subtract to see through it does not include this wait. [#1668](https://github.com/mudler/vllm.cpp/issues/1668) item 4 landed `instrument_seconds` precisely so the residue could be split into "the render did not name this" and "the instrument spent this"; this wait falls into neither. Found by the fresh review of [PR #1711](https://github.com/mudler/vllm.cpp/pull/1711) as a by-product of attempting the per-site charge gate of [#1718](https://github.com/mudler/vllm.cpp/issues/1718), and it is why that attempt could not be made honest. NOT FIXED IN FLOW: a production change to a shared instrument that four render paths run through, owing its own red-first evidence and its own mutation. It also interacts with the charge clamp #1711 lands -- a `Close` that charges a pre-lock interval charges it to the record about to end, so the per-target high-water mark seeded at `o.start` is what keeps the charge inside `[start, end]` and the conservation invariant intact; doing it before that clamp existed would have broken the invariant again. Owned by [`ltx25-phase-instrument.md`](../specs/ltx25-phase-instrument.md) `## Owed` | bug | -| [#1720](https://github.com/mudler/vllm.cpp/issues/1720) | `LTX25-PHASE-INSTRUMENT` | **`WriteJson` reads `Elapsed()` and `Records()` under TWO separate acquisitions of the process-wide mutex, so `wall_seconds` and the record set are no longer one snapshot.** [PR #1711](https://github.com/mudler/vllm.cpp/pull/1711) moves `PhaseLog::WriteJson`'s clock read ABOVE its `ByStart(Records())` so the writer's own copy and sort stop being charged to the render's wall and therefore to `unaccounted_seconds` -- the defect [#1569](https://github.com/mudler/vllm.cpp/issues/1569) tracks -- and this issue is the cost of that repair recorded rather than hidden. Under the previous order the pair was effectively one snapshot in the direction that matters, because the clock was read LAST and `wall >= max(end_seconds)` held by construction; it no longer does. The observable if it broke is a NEGATIVE tail gap, which `gaps` reports and which `ltx2 phase log: the emitted table DECOMPOSES its residue into the gaps between leaves` refuses at `CHECK(seconds >= 0.0)`, so the table would say so rather than pass quietly. It is UNREACHABLE on the shipped path, and that is a property of the CALL SITE rather than of the function: both `WritePhaseLog` calls in `src/vllm/multimodal/ltx2_video.cpp` run after `generate_span.Close()`, that span is the last live scope, and `PhaseLog::Close` stops and JOINS the sampler before it returns when nothing is left live, so no thread can close a scope between those two lines on any path this project ships. A fresh review also failed to stage the inversion adversarially: 27,471 probes of a churn thread against a replica of the two statements produced zero. NOT FIXED IN FLOW: the real repair is to make the pair a single locked snapshot, and `Elapsed()` and `Records()` are separate public entry points on `PhaseLog`, so a combined one is a public API change and owes its own row, spec and red-first evidence rather than being smuggled into a repair commit for a different defect. Owned by [`ltx25-phase-instrument.md`](../specs/ltx25-phase-instrument.md) `## Owed` | bug | -| [#1743](https://github.com/mudler/vllm.cpp/issues/1743) | `LTX25-DIT-ATTN-FLASH` | **The DiT attention swap on `main` (`90e8c3c85`, [#1549](https://github.com/mudler/vllm.cpp/issues/1549)) changes what LTX-2.5 RENDERS, by more than one pixel of global image shift.** Measured at `768x448/49f`, seed `20260820`, one binary (`834cec55...`), one `dgx:gpu0` lease (`rc` job `acff8e89-d704-4f17-a9f2-d354aba53b0d`), against the criterion registered in `.agents/specs/ltx25-dit-attn-flash.md` §10.4 BEFORE the renders were taken. **All six registered checks fail**: mean \|delta\| RGB **6.414156** against `<= 1.0`, worst-frame PSNR **22.269 dB** against `>= 40` (aggregate 25.822), worst-frame SSIM **0.880694** against `>= 0.99`, V4 luma \|delta\| over adjacent-frame MAD **0.709189** against `<= 0.10`, audio PSNR **29.368 dB** against `>= 40`, audio Pearson r **0.932682** against `>= 0.999`. 0 of 49 frames bit-identical, max \|delta\| 253 of 255, and 98.9-99.7% of the pixels in EVERY frame differ, with a broad unimodal histogram rather than a small mean hiding a tail. **THE CONTROL MAKES IT ATTRIBUTABLE**: a third render, `flash` again on the same binary and seed, is BIT-IDENTICAL (49/49, max \|delta\| 0, PSNR inf, SSIM 1.000000) and passes its own C0 content checks, so `R = 0.000000`, the run-to-run noise floor is nil, and every bit of the delta is the swapped op. Routing is proved two-sided per arm from each arm's own `VT_OP_PROVIDER_STATS=1` log (`op18=0 op21=1` flash, `op18=1 op21=0` naive). **THE KERNEL AGREES WITH ITS REFERENCE WHEREVER THIS TREE CAN MEASURE IT**: `test_ltx2_device` in the same lease reads 22/22 cases, 749/749 assertions, device-vs-host MAXIMUM **8.94e-08** against a committed `2e-5` tolerance (the lines labelled `device-vs-host` run 5.96e-08 to 8.94e-08 and the largest still clears the tolerance by two orders of magnitude, so quote a named family's maximum rather than its smallest reading; the log also carries other families on other tolerances, the bf16 keyframe arm legitimately reading 3.31e-03, so no single number summarises it). The `8 of 8` flash-selection count beside it is a **CPU-backend** routing count on `ReducedParams` - it proves the knob routes at fixture size, NOT that the CUDA kernel is exact at head_dim 128 / 2352 keys / 48 layers, which no gate in this tree reaches; the CUDA routing proof for this run is the render's own op log. `vt::Attention` and `vt::AttentionDenseFlash` differ only in association, and the 4.0e4-1.7e5 single-ULP bf16 flips per forward that §10.2 predicted AMPLIFY instead of damping across the render's **30 sampler steps at 4 DiT forwards each, 120 forwards** - about 2.5 orders of magnitude on mean \|delta\|, and 322x the ±1 LSB dither row. Against one pixel of global horizontal shift (5.183 / 28.1 dB / 0.8705 / 0.624) the delta is worse on **three of four axes**: mean \|d\|, PSNR and V4 are worse, while SSIM is a similarity and the measured 0.880694 worst / 0.901395 mean sit slightly ABOVE the shift's 0.8705. Amplification is the DEDUCTION joining a ~1e-7 op bound to a 6.414 pixel delta, not a measured quantity. Divergence also grows along the frame axis (r=+0.753 index vs mean \|delta\|, r=-0.828 vs SSIM). **CONTEXT, not a control**: the 20260820 baseline (`a50c57d69`, an ancestor, naive path) against today's naive arm reads mean \|delta\| 9.452407, PSNR 22.841 dB, V4 1.026 - so two NAIVE renders across builds diverge MORE than flash-vs-naive at one build, and the trajectory is unstable under any arithmetic perturbation rather than under this one. NOT FIXED IN FLOW and deliberately not: §9 forbids widening a threshold to admit it, no number in §10.4 moved, and whether a **7.112x** arm that renders a different video stays the default is a product decision. Neither arm is established as CORRECT - every figure is a difference between two renders, so an absolute reference is owed and is a different row. Recorded in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) §10.7; evidence `/mnt/nas_share/rc/ltx25-attnflash/pixel-ab/1612-r3/`. Measured under [#1612](https://github.com/mudler/vllm.cpp/issues/1612) | bug | -| [#1723](https://github.com/mudler/vllm.cpp/issues/1723) | `LTX25-DIT-ATTN-FA2-HD128` | **The `docs/USAGE.md` checkpoint-registry row for LTX-2.5's distilled NVFP4 DiT disagrees with the artefact on the share in TWO independent fields, and neither disagreement can be settled from inside this repository.** The row pins `Lightricks/LTX-2.5 @ 6c7e5e573ac1667efc83407806fe9b0b93730e60` and **18,721,548,408 bytes** for `diffusion_models/ltx-2.5-22b-distilled-transformer-nvfp4.safetensors`; that file's own `huggingface_hub` `.metadata` sidecar records `8a4ff96f581e72bedc1b44367581c49d544a05f1`, and `stat -c %s` on the shared copy gives **18,721,432,024 bytes** -- a difference of **116,384 bytes**, far too small to be a different model and far too large to be rounding. Both values were RE-DERIVED rather than transcribed, twice, by two independent sessions. The SIZE is the stronger signal: a size is what this registry uses to identify an artefact when no content hash is available, so a size that differs means the BYTES differ, while a revision can legitimately differ from what a row's author measured. Six `Lightricks/LTX-2.5` sidecars exist on the share and ALL SIX record `8a4ff96f...`, which is what makes that value a snapshot revision rather than a blob id. DELIBERATELY NOT fixed by editing the row: both fields admit the same benign explanation -- a later re-quantization published under an unchanged name, re-downloaded after the row was written -- and replacing a possibly-stale pin and size with values whose only provenance is "what happens to be on the share today" swaps a possibly-stale pin for a DEFINITELY UNVERIFIED one, which is worse; the two bf16 DiT rows have no sidecar at all, so nothing local contradicts them either. Settling it needs an authenticated fetch at a named revision and a re-hash -- authority no row has held and something no checker here can gate, and an unauthenticated HuggingFace tree API has returned a fabricated content hash for a gated repo in this project's history, so an etag nothing re-derived locally is not a pin. SCOPE: no shipped model arm runs this file today; `LTX25-DIT-ATTN-FA2-HD128` ([#1551](https://github.com/mudler/vllm.cpp/issues/1551)) is gated against the bf16 DiT, the two VAEs and the NVFP4 Gemma-4 encoder, whose sha256s were confirmed by hashing local bytes. SPLIT OUT FROM [#1702](https://github.com/mudler/vllm.cpp/issues/1702) rather than left on it: #1702's subject is the three LTX-2.5 artefacts with no registry row, that half IS fully fixed by #1551's pull request and closes with it, and this discrepancy is a different bug about a different, PRE-EXISTING row that was found while fixing it -- so tracking it on #1702 would have closed the tracking for debt that remains open. Under `## Owed` in [`ltx25-dit-attn-fa2-hd128.md`](../specs/ltx25-dit-attn-fa2-hd128.md) | bug | -| [#1724](https://github.com/mudler/vllm.cpp/issues/1724) | `LTX25-DIT-ATTN-FA2-HD128` | **`LaunchDenseFA2Bf16`'s head-dim `else throw` is unreachable while the admissibility guard above it stands, so no test enters it and no mutation can turn it red.** `cuda_flash_attn_fa2.cu::LaunchDenseFA2Bf16` now ends its dispatch `if (d == 64) { ... } else if (d == 128) { ... } else { throw ... }`. The arm is RIGHT and the fresh review of [PR #1704](https://github.com/mudler/vllm.cpp/pull/1704) asked for it: the head dim is a TEMPLATE parameter, so the two compiled instantiations are two call sites and not one call with an argument, and the shape it replaced put the 128 call in a bare `else` on the reasoning that a head dim with no instantiation would be a LINK error. That reasoning is wrong -- the set the function serves is decided by the `d != 64 && d != 128` guard earlier in the same function, not by the linker -- so widening that guard alone, the exact first edit a head_dim-192 rung makes, links fine and sends 192 into the 128 kernel, which reads 128 of its 192 channels and returns a SILENTLY TRUNCATED answer. The two guards are ORDERED, which is what leaves the arm without coverage: no input reaches it, no test can enter it without first making the very edit it guards, and deleting the arm or inverting its condition leaves every case green. NOT a regression and NOT a capability that landed dead -- the same input previously produced a silently truncated answer with no diagnostic, so the arm strictly replaces silence with a named refusal. The owning suite `test_ops_attention_dense_fa2` is CUDA-gated in all 12 of its cases -- each returns early on `!HasCuda()` -- so doctest reports 12 cases and ZERO assertions on a CPU build, so it gives the arm no coverage on the authoring host either. OWES: the first commit that widens the `d != 64 && d != 128` guard -- for head_dim 192, for f32, or for any new rung -- owes a RED-FIRST case proving the throw fires for a head dim the widened guard admits and the launcher has no instantiation for, taken BEFORE the instantiation that makes the throw unreachable again is added; that widening is the only moment at which the guarantee is both reachable and provable. NOT fixed in flow, because the only change that makes it testable is the guard widening this issue conditions, which is a separate rung with its own row. Under `## Owed` in [`ltx25-dit-attn-fa2-hd128.md`](../specs/ltx25-dit-attn-fa2-hd128.md) | bug | -| [#1734](https://github.com/mudler/vllm.cpp/issues/1734) | `LTX25-DIT-ATTN-FA2-HD128` | **`scripts/ltx25-dit-attn-fa2-hd128-ab.sh` prints an EMPTY `memavail low-water:` for every arm, and the cause is the writer rather than the reducer that prints it.** Found by this row's own A/B on `dgx:gpu0` (rc job `91e0b5d9-b7f7-4b69-bf3f-d593aa25f871`, artifacts `/mnt/nas_share/rc/ltx25-fa2hd128/out/20260822T203535Z/`). At line 367, `n=$(grep -c 'last=' "$log" 2>/dev/null \|\| echo 0)` emits TWO lines when the count is zero, because `grep -c` prints `0` AND exits 1, so the `\|\| echo 0` arm fires as well. Line 369 then writes that value into a tab-separated record, which lands on disk split across two lines: `watch-flash.tsv` is 85 lines with `NF=3`, 85 with `NF=2` and 16 with `NF=4` (186 total), and `watch-fa2.tsv` reads 85 / 85 / 6. The 85 pairs are exactly the 5-second polls taken during the ~7 minutes of model load, before any `last=` line existed; once the log has a match `grep -c` exits 0 and the record is well formed, which is why a minority of lines are intact. TWO consequences. The reducer at line 390 reads `$4` over a file where `$4` is empty on 170 of 186 lines, and an empty string sorts first under `sort -n`, so `head -1` returns it -- that is the reported symptom, and the reducer is independently fragile because it strips the `memavail_gib=` prefix off `$4` alone and so depends on a field position it cannot rely on. The sample cap's own test `[ "${n:-0}" -ge "$WANT_SAMPLES" ]` also receives the two-line value for those 85 iterations; harmless here because the cap cannot fire before a sample exists, but a non-integer reaching an integer comparison inside the guard that stops a job on a shared box is not something to leave standing. TOUCHES NO NUMBER in the measurement: both arms report `stopped_by=sample-cap`, which is the direct evidence that neither was stopped by memory pressure, and re-derived with a prefix-stripping match instead of a positional one the low-water is **40.3 GiB on both arms** against `MEM_FLOOR_GIB=12.0`, so the run stayed 3.36x above its own floor. NOT fixed in the record commit that found it, because that commit is `.agents/`-only by scope and a `scripts/` edit owes a red-first case over a fixture `watch-*.tsv` asserting the record is one line and the low-water reduces to the expected value. Under `## Owed` in [`ltx25-dit-attn-fa2-hd128.md`](../specs/ltx25-dit-attn-fa2-hd128.md) | bug | -| [#1649](https://github.com/mudler/vllm.cpp/issues/1649) | `GATE-WINDOWS-PORTABILITY-TARGET-SCOPE` | **`check-windows-portability.py` read `/w` anywhere in `CMakeLists.txt` as a project-wide negation of `/W4 /WX`, so the vendored BoringSSL targets' PRIVATE `/w` red `windows-msvc-cpu` and `windows-msvc-vulkan` on main and on every PR.** `check()` set `warnings = cmake` -- the whole file -- concatenated `cmake/CompilerWarnings.cmake` and did a flat token search, which cannot tell a global `add_compile_options(/w)` from a `target_compile_options( PRIVATE /w)`; only the first negates the policy. The refusal lands before any translation unit is read, so the job carries no `error C####`. SECOND red, which #1649 did not record: the same defect fails this checker's own suite -- `test_real_tree_msvc_warning_policy_reaches_the_cxx_compile` with `negation='/w'` on main at `8540a2755` (78 tests, 1 failure) -- so the tree asserted the contradiction in two places at once and one fix clears both. FIXED IN FLOW: `without_foreign_target_compile_options` blanks `target_compile_options` spans whose target PROVABLY names only targets this project never declares (`add_library`/`add_executable` first arguments, with `foreach` bindings resolved); anything unresolved stays in scope, and `cmake/CompilerWarnings.cmake` is kept whole because it applies the policy through an unresolvable function parameter. The widening is bounded by three guard properties proved discriminating by MUTATION -- forcing `_target_is_foreign` to `return True` reds exactly those three and nothing else (82 tests, 3 failures), tree restored byte-for-byte. 82 tests OK; the real tree now prints `Windows portability contract OK`. RESIDUAL, stated not hidden: a negation reaching a project target through a `set()` binding rather than a `foreach()` one is still not caught; no such construct exists in the tree today. Recorded under [`gate-windows-portability-target-scope.md`](../specs/gate-windows-portability-target-scope.md) `## 4. Risks / decisions` | bug | -| [#961](https://github.com/mudler/vllm.cpp/issues/961) | `BACKEND-CUDA-SM110` | `tests/tools/test_online_gate_startup.py:259` guards `shellcheck` absence with a check that cannot fire, so an absent instrument reads as a code verdict: `test_serve_low_tools` raises `FileNotFoundError: 'shellcheck'` instead of skipping. Filed by the sm_110 baseline lane ([#955](https://github.com/mudler/vllm.cpp/issues/955)), where the leased `thor:gpu0` worker carries no `shellcheck`, so the baseline names the failure as a known entry rather than an sm_110 fact. Indexed late: the issue was opened 2026-08-15 and its index row was lost with the unmerged repair of PR [#956](https://github.com/mudler/vllm.cpp/pull/956) | bug | -| [#962](https://github.com/mudler/vllm.cpp/issues/962) | `BACKEND-CUDA-SM110` | NVFP4 marlin disagrees with itself on sm_110: `test_ops_moe_grouped:1262` fails `CHECK(bitdiff == 0)` on `NVFP4 block8-vs-block16 M=8 K=4096 N=4096 bitdiff=15/32768`, against a kernel configure reports as `marlin-nvfp4: ENABLED for [110]`. A live kernel defect on an ENABLED feature, not an absent one, and the only substantive standing sm_110 finding in the [#955](https://github.com/mudler/vllm.cpp/issues/955) baseline. Indexed late, same cause as the row above | bug | -| [#1363](https://github.com/mudler/vllm.cpp/issues/1363) | `BACKEND-CUDA-SM110` | Thor reports 30 GiB of swap, all free, measured inside `rc run -d thor:gpu0`, so the "zero swap" half of the box's `THIS BOX REBOOTS INSTEAD OF OOM-KILLING` warning is stale. `vm.overcommit_memory=1` is unchanged and the three 2026-08-11 reboots were observed, so the hazard stands and is not relaxed; unresolved are whether the swap is the host's or a container view, whether it changes the failure mode at all, and when it appeared | bug | -| [#1721](https://github.com/mudler/vllm.cpp/issues/1721) | `ENV-GPU-LEASE-METHODOLOGY` | `git stash` is repo-global across linked worktrees, so a bare `git stash pop` in one worktree consumes whatever sits at `stash@{0}` — which on this checkout belonged to another row. `git stash` on an already-clean tree saves nothing and prints nothing, but the paired `pop` still fires: it pulled `row/ENG-PUBLIC-DOC-PROJECTIONS`'s 56-file entry into an unrelated worktree, and only survived because the conflict made git KEEP the entry. 14 entries are on the stack, some labelled recovery. Same shape as #777 and #998: a resource that looks per-worktree and is per-repository | bug | -| [#1725](https://github.com/mudler/vllm.cpp/issues/1725) | `BACKEND-CUDA-SM110` | `kMatmulFp8Cutlass` and `kMatmulFp8BlockScaled` fall through to the portable CPU tier and SEGFAULT on sm_110 (`test_ops_fp8_cutlass:191`, `test_ops_matmul_fp8_block_cuda:345`, measured at `0764ded2b`). The sm_110 baseline attributed them to [#960](https://github.com/mudler/vllm.cpp/issues/960), which was CLOSED COMPLETED three days earlier by `d607fec4c` -- that fix covered `QuantFp8Static` only, while these two ops are registered from TUs `CMakeLists.txt:1790-1791` builds solely for `VT_CUTLASS_FP8_ARCHS`. `cuda_matmul_fp8_block_cutlass.cu:56-58` asserts they refuse by name instead, which the measurement contradicts; `cffe59b02` has since rewritten that dispatch, and no CI lane can see either way because `cutlass-fp8` is ENABLED on the GB10 gate host | bug | -| [#1731](https://github.com/mudler/vllm.cpp/issues/1731) | `ENG-RECORD-CONFLICT-SURFACES` | `check-agent-record` is RED on `main`: `#1649` is listed twice in `.agents/issue-index.md`, at `:592` (added by `a7bb3130b`, the lane that FILED it) and `:632` (added by `2f2a70925`, the lane that FIXED it). Neither is wrong and neither could see the other -- `merge=union` combines two appends silently, so a duplicate is the ordinary outcome when filing and fixing happen on different branches, and the checker only notices once both have landed on main, where every later PR inherits the red. Not repaired in flow: the index preamble forbids editing or deleting a row, and the two bodies carry different facts (`:632` has a second red and the mutation evidence, `:592` has the attribution to `a50c57d69`), so choosing which survives is a judgement | bug | -| [#1358](https://github.com/mudler/vllm.cpp/issues/1358) | `ENG-MM-INPUT-PIPELINE` | `LoadQwen3VLWeights` reads the whole vision tower into `Qwen3VLWeights::vision` (`src/vllm/model_executor/models/qwen3_vl.cpp:418`, on the production path via `qwen3_vl_registry.cpp:97` -> `ModelRegistry::Load` -> `LoadedEngine::FromModelDir`) and NOTHING in `src/` ever reads it back. The only consumers of `Qwen3VLWeights` anywhere are three hardware e2e tests (`test_qwen3vl_e2e.cpp:110`, `test_qwen3vl_video_e2e.cpp:152`, `test_qwen3vl_registry_e2e.cpp:120`); the registered forward consumes ALREADY-MERGED embeddings off `ModelForwardInput.mm` and never touches `vl.weights().vision`. So on the server path the tower is paid for at load — widened bf16 -> host f32 — and never used. This is `.agents/reachability.md`'s unpassed-parameter shape wearing a loader's clothes, and unlike the usual case it costs memory rather than only being dead. FOUND while enumerating every production tower-load site for #607 L3, and it is why the L3 RSS measurement uses Muse Glimmer rather than this site. NOT fixed in flow: wiring the tower into the server's mm forward is the MM-SERVE-E2E residual `server_main.cpp:1315-1321` already names, a feature with its own spec and gate rather than a repair. What L3 does add is the flag that stops paying for it — `--language-model-only` now skips this exact load. Owned by `ENG-MM-INPUT-PIPELINE`; listed under `## Owed` in [`multimodal-track.md`](../specs/multimodal-track.md) §1.5 L3 | bug | -| [#1359](https://github.com/mudler/vllm.cpp/issues/1359) | `ENG-MM-INPUT-PIPELINE` | Muse Glimmer's perception encoder is held in HOST F32 where the checkpoint ships bf16: `MuseGlimmerVisionWeights`, `MuseGlimmerVisionAdapterWeights` and `MuseGlimmerVisionTower::projection` are all `std::vector` (`include/vllm/model_executor/models/muse_glimmer_vision.h:106-118`), so the 809 vision tensors that are 3.580 GiB on disk cost **7.161 GiB resident** — twice the checkpoint's own — with no annotation naming a reason, against AGENTS.md's "Inherit vLLM defaults". Measured from the two shard headers of `/mnt/nas_share/checkpoints/muse-glimmer-30b` (1436 tensors, 55.463 GiB total; the encoder is 6.45% of it). Nobody noticed for the reason that paragraph names: a token gate cannot detect a dtype that is too wide, so `test_muse_glimmer_vision` and the wiring gate are both correct and both blind. Qwen3-VL's tower has the same shape (`qwen3_vl_vision.h:76-82`), and `qwen3_vl_vision.h:106-114` already records the host-f32 form dominating encode time without recording that it is also a polarity departure. NOT fixed in flow: narrowing a tower's storage dtype changes numerics on every path that reads it, so it needs its own spec, a bf16-vs-f32 comparison against the reference and its own gate — the surprising-fix path, not an in-flow repair — and it is orthogonal to #607 L3, which removes the tower entirely at zero limits rather than narrowing it. Recorded in [`multimodal-track.md`](../specs/multimodal-track.md) §1.5 L3 beside the RSS threshold, which is stated against 7.161 GiB for this reason. Owned by `ENG-MM-INPUT-PIPELINE`; listed under `## Owed` | bug | -| [#1733](https://github.com/mudler/vllm.cpp/issues/1733) | `ENG-RECORD-CONFLICT-SURFACES` | CLOSED as a duplicate of [#1731](https://github.com/mudler/vllm.cpp/issues/1731), filed seventeen minutes after it against the same red, and recorded here rather than left unindexed because its one MEASURED claim is FALSE and an append-only row cannot be corrected in place later. Both issues report that `.agents/issue-index.md` lists [#1649](https://github.com/mudler/vllm.cpp/issues/1649) twice on `main` at `038ff61e5`, at `:592` from `a7bb3130b` (the lane that FILED it) and at `:632` from `2f2a70925` (the lane that FIXED it). #1731 is the earlier filing, already carries an index row, and is the record. #1733 adds one thing #1731 does not, and it is wrong: "the repair is measured and it is NOT blocked", on the evidence that removing one of the two rows in a worktree and running both checkers gives `agent record OK: ENGINE=170 MODEL=377 ...` and `OK: issue index append-only`. The second half is an artefact of the instrument. `scripts/check-issue-index-append-only.py:50-51` diffs `merge-base(origin/main, HEAD)..HEAD`, which reads COMMITS, so an UNCOMMITTED deletion is invisible to it. Measured on `row/FIX-ISSUE-INDEX-1649-DUP` at base `038ff61e5`: deleting `:592` in the WORKING TREE alone returns `OK: issue index append-only` at rc 0 with `git diff --numstat 038ff61e5..HEAD -- .agents/issue-index.md` EMPTY, and committing the byte-identical deletion turns the same checker rc 1 with a `removed:` line naming the row. #1733's own quoted `agent record OK: ENGINE=170 MODEL=377` is the tell, because that is the working-tree reading and the committed tree cannot produce it while the duplicate stands. So the duplicate IS base-reachable, the two checkers ARE in genuine contradiction on this tree, and the repair is the argued exception #1731's row anticipated rather than the free edit #1733 reported | bug | -| [#1755](https://github.com/mudler/vllm.cpp/issues/1755) | `LTX25-PHASE-INSTRUMENT` | `PhaseLog::RenderText` read `Elapsed()` AFTER `ByStart(Records())`, so the CONSOLE copy of a phase table charged its own copy and its own sort to the `WALL` it printed and to the `unaccounted` row above it. That is [#1569](https://github.com/mudler/vllm.cpp/issues/1569)'s defect on #1569's own sibling emitter: the file copy was repaired and the copy a reader watching a terminal gets was not. The call site was the larger half -- the console block stood at the END of `WriteJson`, after the whole `nlohmann` object was assembled, so it absorbed the JSON build as well: over five `WriteJson` calls on the 8001-record unit timeline `sum(leaf)` held at 0.189 s while the console's `unaccounted` climbed 0.065 -> 0.134 -> 0.200 -> 0.265 -> 0.329 s, about 66 ms of writer work per call charged to a render that had not run. NOTHING COULD SEE IT: `RenderText` prints every total with `%10.3f` and a copy and a sort of 8000 records is 0.12 ms, a quarter of one step of that format, so applying #1569's own one-line repair to its sibling left the suite at `7 \| 7 passed` and `100 \| 100 passed`. Found by the fresh review of [PR #1711](https://github.com/mudler/vllm.cpp/pull/1711) and fixed in the same flow: the clock is read first in both emitters, the console block moves above the copy, the sort and the build, and the gate takes its bound from `%10.3f`'s own last digit rather than from a wall-clock ratio, over a table large enough for the defect to cross it | bug | -| [#1760](https://github.com/mudler/vllm.cpp/issues/1760) | `LTX25-PHASE-INSTRUMENT` | The `### 10` console call-site gate holds only the JSON BUILD, not the copy and the sort. `M-SITE-MID` slides the `StderrEnabled()` block BELOW `ByStart(Records())` and `Sum(...)` while leaving it ABOVE the `nlohmann` build -- [#1755](https://github.com/mudler/vllm.cpp/issues/1755)'s own class, the console `WALL` charged with the writer's copy and sort -- and it survived **19 of 20 runs** at `test cases: 8 \| 8 passed` and `assertions: 120 \| 120 passed`, with one run in twenty red at `8 \| 7 passed`. A mutation a gate catches once in twenty is one the gate does not catch. The cause is arm (A)'s table SIZE and not its bound: at 16000 records the writer's per-record `phases` build measures 6.5388e-3 to 7.2559e-3 s = 6.54 to 7.26 steps of `%10.3f`, while the copy plus the sort at that same size measures 3.8975e-4 to 1.0344e-3 s = 0.39 to 1.03 steps and straddles the one-step bound. Enlarging the table is measured shut -- `WriteJson` holds ~3.3 KB of `nlohmann` per record while it dumps, so a table big enough costs ~1 GB of resident set through an arm that already costs 141 MB and 2.2 s in CI -- and a new wall-clock tolerance is forbidden by [#1668](https://github.com/mudler/vllm.cpp/issues/1668). NOT [#1718](https://github.com/mudler/vllm.cpp/issues/1718), which is the `instrument_seconds` charge-site class. What would settle it is a bound that does not go through the printed format: a structural assertion over the block's position, or a `RenderText` HANDED the wall it prints. Found by the fresh review of [PR #1711](https://github.com/mudler/vllm.cpp/pull/1711) and reproduced by the session that recorded it | bug | -| [#1736](https://github.com/mudler/vllm.cpp/issues/1736) | `ENG-EXPERT-STREAM-DEVICE` | **On `Qwen3.8-2.4T-A95B UD-Q1_0` the `--device cuda` arm and the `--device cpu` arm emit different tokens, and nothing measures whether CUDA is WORSE or only DIFFERENT.** W0g excluded the router gate weights, the embedding table and the W0f host alias, and named no cause. Three things keep the question open: the CUDA continuation degenerates into a mechanical recursion after the 8 tokens the arms share, which a coin flip between two equally good tokens does not produce; every comparison so far is arm-against-arm with no oracle, so "they differ" cannot say which arm is wrong; and the growth-rate argument for "partly systematic" does not survive a proper fit -- a least-squares fit of `log(divergence)` on `log(block + 1)` over all eight recorded points gives an exponent of 0.651 +/- 0.066, interval [0.489, 0.813], which INCLUDES the 0.5 a random walk predicts and EXCLUDES the 1.0 a systematic error predicts, so the two-point 24x-against-9.5x reading overstated it. Scoped as wave **W0h** of `ENG-EXPERT-STREAM-DEVICE`: feed BOTH arms the identical token sequence through the ABI logits processor (`include/vllm.h` v8, applied at `src/vllm/v1/sample/sampler.cpp:441`) and measure the negative log likelihood each assigns to held-out text, which is a quality statement needing no oracle. The decision rule is PRE-REGISTERED before any measurement, the oracle arm is `llama-cpp-unsloth` at `36fe8e1cc` (`gateable = no`, owed by [#933](https://github.com/mudler/vllm.cpp/issues/933)), and the in-tree `VT_CPU_REF=1` switch is excluded on arithmetic (every tensor to `kExpandBf16` at `gguf_keep_quant.cpp:157` is 4.37 to 4.87 TiB against a 119.631 GiB box). No product code, no speed claim, G0-SPEED stays VOID. Spec [`cuda-arm-degradation-experiment.md`](../specs/cuda-arm-degradation-experiment.md) | verification | -| [#1746](https://github.com/mudler/vllm.cpp/issues/1746) | `V1-LOGITSPROC-HOST-ADDRESSABLE` | **`apply_logits_processors` hands every ABI logits-processor callback a raw `cudaMalloc` pointer on CUDA/GB10, because it gates the staging bounce on `Backend::UnifiedMemory()` where the question is `Backend::DeviceMemoryIsHostAddressable()`.** `src/vllm/v1/sample/logits_processor/builtin.cpp` sets `host = logits.data` when the WIDE predicate holds, and `CudaBackend` answers it `pageable_memory_access && integrated`, which is true on GB10 over allocations `CudaBackend::Alloc` takes from `cudaMalloc`; CUDA never overrides the narrow predicate, so it keeps the base `false` from `include/vt/backend.h`, whose own comment says a backend must opt in "because being wrong here hands a device pointer to a host memcpy and segfaults". This is the [#844](https://github.com/mudler/vllm.cpp/issues/844) / [#1435](https://github.com/mudler/vllm.cpp/issues/1435) / [#960](https://github.com/mudler/vllm.cpp/issues/960) class in a second location, and `src/vt/op_provider.cpp` warns about it in the same tree. The CPU suite could not see it: on `Device{kCPU,0}` both predicates are true and the pointer really is host memory, so the wrong one reads correct. FIXED IN FLOW: the predicate narrows, the `else` staging arm is unchanged, and a new own-executable test carries the GB10 pair on a fake backend — `UnifiedMemory()` true, `DeviceMemoryIsHostAddressable()` false — entering through `Sampler::forward` rather than through the function, so the reachability mutation of the one production call site turns it red. A second case pins that a backend answering both predicates true keeps the zero-copy in-place wrap, so the fix does not become "always stage". Spec [`logits-processor-host-addressable.md`](../specs/logits-processor-host-addressable.md) | bug | -| [#1748](https://github.com/mudler/vllm.cpp/issues/1748) | — | **`CpuBackend` never opts in to `DeviceMemoryIsHostAddressable()`, so every reader of that predicate gets the conservative `false` on the one device where it is trivially true.** `src/vt/cpu/cpu_backend.cpp` overrides `UnifiedMemory()` to `true` and leaves the narrow predicate at the `include/vt/backend.h` default, although `CpuBackend::Alloc` returns ordinary aligned host memory, `Copy` is `std::memcpy`, and the class comment two lines above says "Host and device memory are the SAME allocation here". Found while grounding [#1746](https://github.com/mudler/vllm.cpp/issues/1746). Three readers: `ReferenceTierEligible` and `ReferenceTierRefusalReason` (`src/vt/op_provider.cpp`) never reach it, because both return earlier on `device == DeviceType::kCPU` — which is why the wrong answer stayed invisible; the direct-upload adoption in `src/vllm/model_executor/models/qwen3_5_weights.cpp` (two sites) returns early on it and therefore never adopts on CPU, and whether adoption is even meaningful there is NOT established and is part of what this issue owes; and `apply_logits_processors` reads it from #1746 onward, so CPU takes a staging bounce of `[n, vocab]` f32 down and back per step, charged only to a request that registered a processor. NOT fixed in flow: the one-line override flips the qwen3_5 residency path and needs its own red-before test and its own measurement, and a crash-class correctness repair must not carry an unmeasured residency change. Owed under [`logits-processor-host-addressable.md`](../specs/logits-processor-host-addressable.md) `## Owed` | bug | -| [#1751](https://github.com/mudler/vllm.cpp/issues/1751) | `LTX25-DIT-ATTN-FA2-HD128` | **`VLLM_LTX2_DIT_FLASH_ATTN` matches the naive arm on a PREFIX and falls through to the FA-2 default for every value it does not recognise, so a typo cannot refuse.** In `src/vllm/model_executor/models/ltx2_device.cpp` the DiT self-attention dispatch reads the three-way A/B knob with two different rules: the naive arm tests `arm[0] == '0'`, so `0x`, `07`, `0flash` and any other string starting with `0` all select `vt::Attention`, while the flash arm uses `strcmp(arm, "flash")` and matches only the exact value. Everything that matches neither -- a typo such as `falsh`, a case difference such as `FLASH`, a trailing space, `naive`, `1`, or an empty string -- lands in the bare `else` and runs `vt::AttentionDenseFa2`, the default, with no diagnostic. WHY IT IS A CORRECTNESS PROBLEM: `flash` is the DENOMINATOR of the 2.74x ratio recorded in [`ltx25-dit-attn-fa2-hd128.md`](../specs/ltx25-dit-attn-fa2-hd128.md) section 8, so a mistyped denominator arm does not fail -- it runs the numerator's kernel a second time and yields ~1.00x, which is also exactly what "no speedup" looks like, so the number cannot report its own failure. That is the same shape `847e22f80` already repaired once on this knob, where the `flash` arm was protected by nothing executable. NOT EXPOSED IN THE RECORDED MEASUREMENTS: `assert_arm_op` in `scripts/ltx25-dit-attn-fa2-hd128-ab.sh` reads the `VT_OP_PROVIDER_STATS=1` announcement and exits 47 on a mismatch, so both arms of section 8.2 asserted their rung rather than assuming it. A human operator who sets the environment variable by hand, in a shell or a service unit, gets no such check and sees a run that looks entirely normal. THE FIX is an exact three-way parse -- `strcmp` against `0` and `flash`, unset for the default -- with a refusal that names the variable and lists what it accepts, in the shape AGENTS.md requires of an unimplemented arm. NOT FIXED IN FLOW, because a product edit to a dispatch that decides which kernel runs owes a RED-FIRST case asserting the refusal fires for an unrecognised value, taken before the parse changes; that is a separate rung with its own row, and the change that found this one is `.agents/`-only by scope. Found by the fresh review of [PR #1704](https://github.com/mudler/vllm.cpp/pull/1704). Under `## Owed` in [`ltx25-dit-attn-fa2-hd128.md`](../specs/ltx25-dit-attn-fa2-hd128.md) | bug | -| [#1727](https://github.com/mudler/vllm.cpp/issues/1727) | `LTX25-DEVICE-RESIDENCY` | **CONSTRAINT ON ANY RE-PROPOSAL of a sampler-based span-slack denominator, which is NOT on `main`: the boundary sampler carries three terms and a gate can falsify only one.** The estimator this row built was measured and WITHDRAWN -- see [`ltx25-device-residency.md`](../specs/ltx25-device-residency.md) `### The span-slack bound, FIFTH and SIXTH shapes` -- so what follows is what anyone rebuilding it already owes. `InstrumentCeiling::Loop` timed the process-wide mutex, `SampleNow()`'s `/proc/self/statm` read, and one flushed write standing for `EmitLocked`. A fresh review of [PR #1712](https://github.com/mudler/vllm.cpp/pull/1712) deleted the flushed write entirely -- the term the derivation argues corresponds to the live lane's per-boundary line -- and the case stayed GREEN 3 of 3 at `compile_status=0`, `assertions: 606 \| 606 passed \| 0 failed`, with the ceiling reading 3.24-9.19 ms against an unmutated range of 0.94-17.5 ms: nothing distinguishes the term's removal from run-to-run noise. That same measurement is what freed the write to be pointed at `/dev/null` instead of fd 2, where at 1 kHz it corrupted 27% of the `[render]` lines [#1413](https://github.com/mudler/vllm.cpp/issues/1413) exists for (158 occurrences in one capture, 116 beginning a line) -- **but that redirect went out with the withdrawn shape and is NOT in the tree**, on `origin/main` or on this branch, so #1413's stream corruption is UNFIXED and nothing here may be banked against it. What stands is the measurement, not a repair: the sink is not what the term measures. TWO further halves of the estimator are UNMEASURED rather than ungated. The SANITIZER lane: the third shape of this bound needed its own constant under ASan and TSan because a sanitizer instruments the scope boundary this quantity IS, and the sampler now inherits that automatically by compiling into the same binary -- but the sample count grows with the slowdown too, so the maximum is drawn from a deeper quantile, and nobody has run `sanitize-cpu` against the estimator (the review could not: a second full build tree does not fit at 93% disk). The FULL-SUITE regime: [#1576](https://github.com/mudler/vllm.cpp/issues/1576) measured this quantity at 3.223 ms inside the full 102-case `ctest` run against 18.8 us under a `-tc=` filter at comparable load, **171x**, and that swing is the whole reason the constant was declared unsafe -- yet both the row's 120-run validation and the review's 24-run one were taken almost entirely under `-tc=`. One full-binary run (containment case as #40 of 102) was indistinguishable from the filtered ones at ceiling 2.31 ms, bound 9.22 ms, worst slack 86 us, so the 171x did not reproduce at n=1. NOT FIXED IN FLOW: gating a term of the estimator is a second instrument over the first, and each unmeasured regime needs a build tree this box cannot hold beside the one it has. What would settle it: one `sanitize-cpu` build with the containment case run 40+ times, and one 40-run population through the FULL binary, both reported as distributions. Owned by [`ltx25-device-residency.md`](../specs/ltx25-device-residency.md) `## Owed` | verification | -| [#1728](https://github.com/mudler/vllm.cpp/issues/1728) | `LTX25-DEVICE-RESIDENCY` | **HYPOTHESIS, filed unmeasured and labelled as such, and a CONSTRAINT ON ANY RE-PROPOSAL rather than a live defect, because the sampler it is about was measured and WITHDRAWN and is not on `main`: on a ROCm runner a 1 kHz boundary sampler inside `(1c)` would call `hipMemGetInfo` 1000 times a second under the process-wide phase mutex.** Raised by the fresh review of [PR #1712](https://github.com/mudler/vllm.cpp/pull/1712) on a CPU-only box with no GPU lease taken. `InstrumentCeiling` runs a thread calling `vllm::multimodal::phase::SampleNow()` once per millisecond for the whole of `ltx2 video: the three carrying phases contain their work and the load keeps its order` -- 3,764 to 27,975 draws per run, measured. `SampleNow()` takes the phase mutex and runs `SampleLocked()`, which invokes the installed `DeviceByteProbe`; `FixtureParams` sets `mp.device = 0` and `src/vllm/multimodal/ltx2_video.cpp` installs a probe calling `vt::Backend::DeviceMemoryInfo`. On the lanes the review could run this is harmless because `CudaBackend` does not override it and returns the -1 sentinel ([#1126](https://github.com/mudler/vllm.cpp/issues/1126)), but `src/vt/rocm/rocm_backend.hip:373` DOES override it. If the reading is right the consequence runs in two directions at once: the render slows because it contends for that mutex against a thread holding it across a driver round trip, and the sampler's own draws inflate, widening the very bound (1c) asserts against. What would settle it, cheaply: one lease on a ROCm device running that case and comparing the printed whole-case boundary median and maximum against this box's 18.4-22.9 us / 0.94-17.5 ms, plus the case wall. If it reproduces, the repair is a null device probe for the sampler's lifetime, or a device-counter cadence far below the boundary-timing cadence. NOT FIXED IN FLOW: the box that found it has no ROCm device and the session was CPU-only, so both reproduction and repair are unmeasured. Filed rather than left in a review transcript, because a 1 kHz driver call under a shared mutex is the kind of thing discovered as "the ROCm CI is slow" months later. Owned by [`ltx25-device-residency.md`](../specs/ltx25-device-residency.md) `## Owed` | bug | -| [#1696](https://github.com/mudler/vllm.cpp/issues/1696) | `TEST-TT-CAPTURE-DECLINE-AMBIENT` | **The capture-decline arms that landed with #1630 (`333509dc5`) are vacuous under an ambient `VT_TT_HOST_FREE_DECODE=0`: the opt-in CHECK reds 833/834 on the unit gate's opt-out leg while the default leg stays green (why CI never saw it).** `support_static_graph_mode()` is the conjunction `HostFreeDecodeEnabled() && VT_TT_DECODE_CAPTURE` (`src/vllm/platforms/tenstorrent.cpp:81-84`), so the 3-arm block's opt-in cell implicitly assumed host-free ON. FIXED IN FLOW: the arms became the full four-cell truth table with BOTH envs pinned per cell and the ambient state saved/restored — mutation-proofed (dropping the capture conjunct reds cell 2, dropping the host-free conjunct reds cell 4) and verified 835/835 on BOTH ambient legs on the P150. The fix missed the #1630 merge window by minutes; re-landed as its own row | bug | -| [#1770](https://github.com/mudler/vllm.cpp/issues/1770) | `VT-CONV1D-TIME-BLOCK` | **The two b0 op-level losses the `out_channels * kernel <= in_len` condition was derived from do not reproduce on a second boot, and the condition measures NEUTRAL on the window.** Found while closing [#1683](https://github.com/mudler/vllm.cpp/issues/1683) in `rc` job `16b594ec-7987-4cae-b377-414adbe0f944` on `thor:gpu0`, worker `rc-worker-kk96r`, boot id `e2112cac-660b-434e-911d-33cbd29b9176`, with arms A, C and D built from `origin/main` at `8eecc05a9` inside one lease and differing only in the row's own files. `.agents/specs/vt-conv1d-time-block.md` §2b prices the UNCONDITIONAL arm C against the baseline at 86 latents and reads `b0_res_conv1` 0.82x and `b0_res_conv2` 0.89x; §3b's condition exists to decline exactly those two shapes and §9 records unconditional blocking as REJECTED on that evidence. The second job, same instrument, same length, three alternated rounds, medians: `b0_res_conv1` C 0.04073 s against D 0.03824 s = **1.065x**, keeping the direction at a quarter of the size, and `b0_res_conv2` C 0.00897 s against D 0.01134 s = **0.791x**, REVERSED -- blocking that shape is 21 % faster and the rule declines it. Over the two shapes together C reads 0.04970 s against D's 0.04958 s, a tie. The window agrees the condition is neutral: at 86 latents the rule decides differently on exactly four shapes, `vocoder.conv1d` makes 54 calls per window so they run 2, 2, 6 and 6 times, and the per-call deltas bound the condition's whole window effect at **0.36 ms, 0.01 % of a 3.5 s window**; the paired 7-round pairing reads C 3.3899 s against D 3.5072 s (0.97x), which at 3 % therefore CANNOT be the condition, arm D's legs being bimodal at 3.394-3.429 and 3.507-3.530 s against arm C's 3.368-3.468 s; at 20 latents the arms are within 1.2 % at every thread count from 1 to 14. Nothing is broken: the shipped arm is correct, bit-identical at two fingerprints across three arms and five thread counts, and scales 11.54x of 14. What is unsupported is that the condition BUYS anything on this box, so either the b0 readings are boot-sensitive at a magnitude the row did not budget for, or one of the two jobs measured something other than the geometry it named. NOT FIXED IN FLOW: settling it needs a fresh lease and a per-geometry spread rather than a median of three rounds, at both 20 and 86 latents, plus the probe's own `--control` residency sweep at the b0 footprint. Owned by [`vt-conv1d-time-block.md`](../specs/vt-conv1d-time-block.md) `## Owed` | verification | -| [#1715](https://github.com/mudler/vllm.cpp/issues/1715) | `BACKEND-TENSTORRENT-GDN` | **The Qwen3.5/3.8 GDN-hybrid family cannot run on Tenstorrent: the GDN linear-attention op chain (`kGdnPrefill`, `kGdnDecode`, `kL2Norm`, `kRmsNormGated`, `kCausalConv1dFwd`/`kCausalConv1dUpdate`, `kGdnStateGather`/`kGdnStateScatter`) has no TT kernel, and the P150 is discrete so an op miss refuses by name rather than falling back** (`src/vt/op_provider.cpp` `Resolve`; `src/vllm/platforms/tenstorrent.cpp:55` allow-lists only OPT/Qwen3-dense/Mistral). Planned at the 2026-08-22 Qwen3.8 gap pass as the family's hard prerequisite on TT. Substrate anchor: the pinned tt-metal ships `ttnn::transformer::chunk_gated_delta_rule` (FLA chunked GDN forward, on-core state), mapping onto `kGdnPrefill` behind a varlen+state-permute adapter; decode composes the rank-1 update from matmul+eltwise with a device shadow keyed by host pointer (the `PagedKvShadow` pattern). Owned by this issue; spec [tenstorrent-gdn.md](../specs/tenstorrent-gdn.md) | feature | -| [#1765](https://github.com/mudler/vllm.cpp/issues/1765) | `GATE-METAL-MLX-COMPILE` | **[#1692](https://github.com/mudler/vllm.cpp/issues/1692) says `src/vt/metal/metal_mlx_provider.mm` is "compiled by NO job in this repository". MEASURED, and that is wrong in a way that changes what is owed:** `.github/workflows/release.yml:347` `mlx_arm64` builds it on `macos-15` against the real `mlx==0.32.0` wheel. The defect is the TRIGGER -- `release.yml` fires on a `v*` tag or a manual dispatch, never on a pull request and never on a push to `main`, and `ci.yml` has no Apple runner (all 17 `runs-on` lines are ubuntu-latest x14, ubuntu-24.04-arm x1, windows-2022 x2; its single `macos` string is `test_release_macos_metadata.py` at `:199`, which compiles nothing). So the only build of the file happens AFTER a change lands, and a break presents as a BLOCKED RELEASE rather than a red check; #1584's edit to `MlxFallback` is the case that proved it. On a configured CPU tree the file has **0** occurrences in `build.ninja` and **0** in `compile_commands.json` -- as does every other `.mm` -- against 3 for `src/vt/op_provider.cpp` as a control. **Closed by a never-linked OBJECT library** (`vllm_metal_mlx_provider_syntax_check`, the `vllm_rocm_platform_syntax_check` idiom) that compiles the file on Linux: it is the ONLY Metal TU that can be, carrying 0 Objective-C constructs against 10/5/19 for `metal_ops.mm`/`metal_backend.mm`/`metal_context.mm`, because `metal_context.h:22` hands it every Metal handle as a `void*` on purpose. `LANGUAGE CXX` for `-x c++`; `-Wno-deprecated` for `#import`, MEASURED narrow (a `[[deprecated]]` call still fails as `-Werror=deprecated-declarations`); guarded `NOT VLLM_CPP_MLX` (the LANGUAGE property is directory-scoped and must not reach the shipping build) and `NOT MSVC` (`#import` there means a type library). **The limit is stated rather than glossed:** the stubs under `src/vt/metal/stubs` are written from this file's call sites, so the gate proves it compiles against the REAL `vt::` seam -- the #1584 defect class exactly -- and is blind to every MLX API change, which `mlx_arm64` alone can see. Red-before/green-after and the `## Owed` residue in [gate-metal-mlx-compile.md](../specs/gate-metal-mlx-compile.md) | bug | -| [#1783](https://github.com/mudler/vllm.cpp/issues/1783) | `ENG-EXPERT-STREAM-DEVICE` | **`.agents/benchmark-record.md`'s W0f entry transcribes the CUDA arm's 32 ids with `7172` dropped in TWO places, so every record that inherited it names the wrong divergence point, and the branch-force run that caught it also FALSIFIES W0h ground 1.** With `7172` restored the two arms share **eight** generated tokens and first diverge at **position 9**, `279` (" the") against `9338` (" France"); the records said six tokens, step 7, and `7172` against `303`. Step 7 is a step both arms AGREE on, which is what the W0f entry's own top-2 table says (`7172` 18.779411 over `303` 18.514702), and the divergent step is step 9 at margin 0.022802. The error is visible with no log at all: [`cuda-arm-degradation-experiment.md`](../specs/cuda-arm-degradation-experiment.md) quotes the CUDA text as "a city **located** in France" and " located" IS `7172`, while the transcribed ids decode to "a city in France". **Inherited at SEVEN sites, enumerated because two earlier passes over this fact each under-counted them**: [`expert-stream-device-slots.md`](../specs/expert-stream-device-slots.md) at FOUR -- its W0e G0-CORRECT bullet, its `## Gates` G0-CORRECT cell, its W0f narrative and its `## Owed` ratified-gate cell; the PUBLIC [`docs/models/qwen3-8-2-4t.md`](../../docs/models/qwen3-8-2-4t.md) at TWO -- the divergence paragraph and the coin-flip clause 33 lines below it, which asserted the argument this same change falsifies; and the SHIPPED HEADER `include/vllm/model_executor/models/qwen3_5_weights.h`, whose `kDeviceAliasAlignment` comment carried the same 0.264709 attribution. The append-only [`benchmark-record.md`](../benchmark-record.md)'s W0g entry then carries TWO further defects rather than one, and they are DIFFERENT defects that have to be searched for separately: an EIGHTH restatement of the wrong margin, and about 200 lines below it the coin-flip inference itself, whose token count is RIGHT ("agree for 8 tokens") and whose ARGUMENT this run falsifies, which is why four sweeps that searched for the number walked past it. The entry keeps its bytes and each defect gains its own INSERTED note rather than an edit. A stale W0e-era sentence in the same spec also still said "Both continuations are coherent" while two other places in that one file said the CUDA continuation degenerates. **The substantive half is ground 1 of W0h** ([#1736](https://github.com/mudler/vllm.cpp/issues/1736), which stays OPEN): it argued the recursion proves the CUDA arm is worse because "a coin flip between two equally good tokens does not produce that". A five-arm CPU-only branch-force run on `dgx:gpu0` (2026-08-23, source `ff8f728071bd5`, one binary, identity arm first and byte-exact against the four-times-recorded CPU answer including its stream counters) shows it does: prefilled with the eight shared tokens plus `9338` the CPU arm recurses into the CUDA continuation, prefilled with `279` it reproduces the CPU tail with 23 ids byte-identical. Independently, reaching the same branch-point context by PREFILL instead of decode flips the top-2 on the CPU arm alone, `9338` 19.962210 over `279` 19.820848 against the reference decode's `279` 19.850554 over `9338` 19.827751. FIXED IN FLOW: benchmark-record is append-only so the W0f entry keeps its bytes and gains an INSERTED correction note plus a new W0h section, the W0g entry keeps its bytes and gains TWO INSERTED notes, the seven inheriting sites listed above are corrected in place and each says what it used to say, and W0h ground 1 is rewritten as OFFERED, TESTED, FALSIFIED rather than deleted. The pre-registered decision rule is byte-identical across the change and the spec carries the one-command digest check that proves it. G0-CORRECT stays FAILING and G0-SPEED stays VOID; grounds 2 and 3 carry W0h | bug | -| [#1777](https://github.com/mudler/vllm.cpp/issues/1777) | `GATE-FP8-NUMERIC-BOUND` | **Layer 2 of [#1189](https://github.com/mudler/vllm.cpp/issues/1189)'s `## Gate design` had no implementation, and nothing in the tree could see a per-tensor scale collapse.** Layer 1 passed on 2026-08-23 (`rc` job `2911ed39`, `TOKEN_VERDICT=PASS`), and the header of `tests/parity/test_qwen27n_fp8_tower_paged_engine.cpp` records why that is not enough: a x1.02 AND a x1.10 perturbation of every FP8 weight scale were demonstrably REACHED on the dgx production build and both still produced 16/16 IDENTICAL tokens, with only x2.00 failing. So a wrong-but-close scale and a silent dequant -- which is numerically BETTER than the quantized path -- are invisible to every token comparison here. The audit corrects the brief in one place rather than re-implementing what exists: the SHAPE ASSERTION at the GEMM boundary is already DONE in two places, `src/vt/ops.cpp::MatmulFp8BlockScaled` asserting `b_scale == [cdiv(N,block_n), cdiv(K,block_k)]` (pinned by `test_ops_matmul_fp8_block_cpu.cpp` G5) and `Fp8BlockScaledRefusalFor` refusing `N % 128` / `K % 128` BY NAME before CUTLASS can answer `Invalid status` (#1453, pinned by G4/G6), and the CUDA TU's own comment at the refusal site says so. FIXED IN FLOW: `tests/vllm/model_executor/models/test_fp8_block_numeric_bound.cpp` bounds every projection against the fixture's independent `double` reference through the production `Fp8BlockLinearMethod::Apply`, TWO-SIDED in the same case -- clean UNDER the bound and x1.10 OVER it, per shape per dtype, plus a 4x margin over the grid -- because a one-directional bound stops biting the moment somebody widens it and nothing would say so. Measured over six shapes: bf16 clean 3.82e-3 against x1.10 1.034e-1 at `kBoundBf16 = 2e-2` (5.2x each way), f32 clean 6.20e-6 against 1.000e-1 at `kBoundF32 = 1e-4`. The statistic floors its denominator at 1% of the projection's own maximum, and that was MEASURED rather than assumed: unfloored, the clean f32 reading moves 835x on one shape because a single output element cancels to near zero, which is a number that cannot carry a bound. `dense_fp8_block::Fp8BlockScaleSpread` is the scale-variance probe #1189 asks for, a max/min RATIO because the issue pins the degenerate reading at exactly 1.0 where a variance reads 0.0, carried with the GEMM counter in one `Fp8BlockStats` snapshot; single-cell grids are counted APART from collapsed ones because 11 of the model fixture's 13 grids hold one cell and a probe that reads 11 on a healthy model is one nobody believes. The split arm also gained the two GEMM-boundary operand assertions the merged arm has carried since M6 -- without the packed one, a short buffer is read PAST ITS END by the first GEMM and the call returns a value. Seven mutations in both directions, tar-snapshotted and `sha256sum -c` verified after each: x1.10 RED 14 assertions, x1.02 RED 14, the same 14 (not predicted: the bf16 readings come in at 2.32e-2 to 2.39e-2 against the 2e-2 bound because the truncating store compounds ON TOP of the scale error, so the bf16 half is caught by only 1.16x and the f32 half by 200x), the bound widened to 1.0 RED 6, a collapsing load path RED 14 across 2 cases, the probe's accounting RED 2, and each new operand assertion RED when defanged. CPU tier throughout, so `build-test-cpu` and `sanitize-cpu` run all of it on every pull request; `cuda-fat-build` runs no `ctest` and nothing here rests on it. **#1189 is NOT closed**: its gate design is now covered but its six milestones' `## Owed` sections are live, and no closing keyword appears in this change. Spec [`gate-fp8-numeric-bound.md`](../specs/gate-fp8-numeric-bound.md) | gap | -| [#1779](https://github.com/mudler/vllm.cpp/issues/1779) | `GATE-PREPUSH-FAIL-LOUD` | **`.githooks/pre-push` named six checkers and three of them had no file, and its file-test guard skipped each missing one in silence while the hook still exited 0** -- so it presented as six gates and ran three, and `core.hooksPath` is set to `.githooks` here, so it runs on every push. `check-policy.py` and `check-state-record.py` went with `0f3e44eee`, `check-public-doc-tables.py` with #1714; all three are deleted, not renamed. PART 1 FIXED IN FLOW: the loop now refuses a name it cannot find, the three dead names are pruned, the dead `--base` case arm goes with them, and `.githooks/README.md` stops listing the retired table gate. A red-first suite executes the hook against a scratch repository and pins both directions. PART 2 IS NOT FIXED AND STAYS OWED under `## Owed` in `.agents/specs/gate-prepush-fail-loud.md`: 65 specs still name a deleted checker, which needs its own row | bug | -| [#1793](https://github.com/mudler/vllm.cpp/issues/1793) | — | The GGUF Qwen3.5/3.6 MoE loader (`qwen3_5_gguf_weights.cpp:1082-1099` @ `5d638b67e`) keeps `in_proj_b`/`in_proj_a` split across three residency routes (kept-quant slice, bf16-expand with optional V-row reorder, `gdn_expand_nk` orientation), so packed GDN decode stays unreached on every GGUF MoE checkpoint after `GDN-MOE-PACKED-BA` closes the safetensors arm. Split out so the first unit carries one byte-exactness argument, not three. Listed under `## Owed` in [gdn-moe-packed-ba.md](../specs/gdn-moe-packed-ba.md) | gap | -| [#1528](https://github.com/mudler/vllm.cpp/issues/1528) | `LTX25-GUIDED-VIDEO` | **Two comments in the LTX-2.5 guided denoiser named `RetakePipeline` as the one upstream caller that sets `force_uncond_pass`, and NO upstream caller sets it at `fd4ded7f`.** `include/vllm/model_executor/models/ltx2_denoisers.h:144-145` and `src/vllm/model_executor/models/ltx2_denoisers.cpp:111-113`, both from `daeff67f2` (row `LTX25-GUIDED-VIDEO`, [#1092](https://github.com/mudler/vllm.cpp/issues/1092), PR [#1102](https://github.com/mudler/vllm.cpp/pull/1102)). Re-derived at the local `Lightricks/LTX-2` checkout whose `git rev-parse HEAD` is `fd4ded7f2d88d3da713abcdd4ad41ecc4a9314ca`, the revision `.agents/specs/ltx25-guided-video.md:6` pins. `git grep -n force_uncond_pass` over the WHOLE tree, not only `*.py`, returns 11 LINES and not one is an assignment by a caller (the filing counted 10 because it restricted the grep to `*.py`; the two counts agree and name different populations): `utils/denoisers.py:74` declares the `_guided_denoise` keyword, `:102-103` are the two reads, `:267,273,297` are `GuidedDenoiser` storing and forwarding it, `:313,319,357` are `FactoryGuidedDenoiser` doing the same, `utils/samplers.py:692` is the substring `force_uncond_pass=True` inside a raised ERROR MESSAGE, and `packages/ltx-pipelines/CLAUDE.md:76` documents it. `git grep -n 'GuidedDenoiser('` finds SEVEN construction sites, three `GuidedDenoiser` and four `FactoryGuidedDenoiser`, which is a stronger statement than the filing's retake-only check: `a2vid_two_stage.py:230`, `keyframe_interpolation.py:232`, `retake.py:305`, `t2a_one_stage.py:155`, `ti2vid_one_stage.py:221`, `ti2vid_two_stages.py:248`, `ti2vid_two_stages_hq.py:271`. None passes the keyword, so all seven take the `force_uncond_pass: bool = False` default. The flag is a CFG++ affordance per `CLAUDE.md:76`: run the uncond pass even when `cfg_scale=1.0`, because CFG++ still needs the uncond prediction for the ordinary differential equation derivative. Upstream ships the plumbing and turns it on nowhere. TWO SMALLER ERRORS RIDE ALONG: the header cited the declaration as `:76`, which is the DOCSTRING and not the declaration at `:74`; and both comments cited `retake.py:305-311`, whose last line is BLANK, the construction being `:305-310`. THE CODE WAS ALREADY RIGHT and nothing about behavior changed: `Ltx2GuidedDenoiseInputs::force_uncond_pass` defaults to `false`, `grep -rn force_uncond_pass src/ include/ tests/` finds two reads and no assignment, and the condition at `ltx2_denoisers.cpp` mirrors `denoisers.py:102-103` operand for operand. This is a documentation defect ONLY, and no wider scope was taken. THE ANCHOR WAS LOAD-BEARING, which is why it is a bug and not a typo: a reader who trusts it concludes retake is missing a flag upstream sets and sets it, and `force_uncond_pass` adds an uncond pass per PRESENT modality, doubling retake's DiT forward count from 8 to 16 and changing its output. The comment invited the exact divergence it appeared to describe. A comment-only change CANNOT be red-first tested, and this row does not pretend otherwise: the evidence is the upstream re-derivation above, cited by `file:line` at the pin, and an independent reviewer re-derived it rather than reading it. FIXED IN FLOW, comment text only. Two specs were checked and deliberately NOT edited: `ltx25-retake.md` never mentions the field, and `ltx25-guided-video.md:198` states the CONDITION without claiming a caller, which is true at the pin | bug | -| [#1807](https://github.com/mudler/vllm.cpp/issues/1807) | `ENV-LEASE-RUNTIME-STAGING` | A leased `rc` job read its gate checkpoint over CIFS from `/workspace` and every job hand-wrote its own existence-check-then-`cp -a` guard, which accepts a truncated shard a killed copy left behind (`/mnt/nas_share/rc/dedup-gate2/build72.sh:118-125`). Developer direction 2026-08-23: copy NAS -> local once, idempotently. FIXED IN FLOW: `scripts/rc-stage-checkpoint.sh` (manifest-defined completeness via `SHA256SUMS`, marker+size fast path that reads no payload, resumable `.part` copy, post-copy verify, refuses a directory with no manifest), `tests/scripts/test_rc_stage_checkpoint.py` (11 hermetic cases, registered in preflight), and the staging paragraph in `.agents/environment.md`. The two 35B gate checkpoints are staged under `/mnt/nas_share/rc/ckpt/` for `GDN-MOE-PACKED-BA` (#1169) | gap | -| [#1815](https://github.com/mudler/vllm.cpp/issues/1815) | `SAMPLE-PROMPT-LOGPROBS` | **`prompt_logprobs` was computed end to end and then DROPPED at the OpenAI serving layer**, so `{"prompt_logprobs": 4}` returned HTTP 200 with the field absent from the body -- the #925 defect class, a key accepted and silently ignored, and the instrumentation blocker [#1775](https://github.com/mudler/vllm.cpp/issues/1775) names ("no production path exposes a logit vector"). Everything below the wire already existed and was gated: parse at `protocol.cpp:311,479`, validation at `input_processor.cpp:152`, the runner source at `prepare_inputs.cpp:223` + `runner.cpp:2196` (#223), the scheduler slice at `scheduler.cpp:1063`, accumulation at `logprobs.cpp:82`, and `RequestOutput.prompt_logprobs` filled by the output processor. `serving_completion.cpp` and `serving_chat.cpp` then never read it, and neither response type had a field to hold it. RED measured over a real socket on `bacb71109`: `key 'prompt_logprobs' not found` on both endpoints. THREE REQUEST VALIDATORS WERE ALSO UNPORTED and each was a wrong ANSWER rather than a missing one: `{"prompt_logprobs":2,"stream":true}` returned 200 AND STREAMED where upstream 400s, `{"prompt_logprobs":-2}` returned 200, and a non-numeric value 400'd with an nlohmann type message instead of upstream's wording. FIXED IN FLOW by mirroring `completion/protocol.py:601` (`CompletionResponseChoice.prompt_logprobs`), `chat_completion/protocol.py:126` (`ChatCompletionResponse.prompt_logprobs`, TOP-LEVEL rather than per choice), `generate/base/serving.py:305-317` (`clamp_prompt_logprobs`), and the two identical `mode="before"` validators at `completion/protocol.py:474-499` / `chat_completion/protocol.py:763-793` as one shared PREFIX plus a per-endpoint SUFFIX, because reading upstream showed the two are NOT interchangeable and a first pass that merged them broke a shipped capability: the completion count has no `-1` sentinel and refuses every negative (`completion/protocol.py:495-499`), while the CHAT count carries `-1` for "every vocabulary entry" (`chat_completion/protocol.py:784-790`) -- which this tree already serves end to end through `ChatTopLogprobs`, gated by `test_serving.cpp` "serving_chat: top_logprobs=-1 returns every vocab entry per token", a case that sets the struct field DIRECTLY and so could not have caught the HTTP surface losing it. The chat suffix also carries `when using \`top_logprobs\`, \`logprobs\` must be set to true` (`:792-796`), which we accepted and ignored. That closes the request-validation half of [#249](https://github.com/mudler/vllm.cpp/issues/249) and the completion-surface divergence [logprobs-all-sentinel.md](../specs/logprobs-all-sentinel.md) records under "Scope". `prompt_logprobs: 0` with `stream` still parses, because upstream's condition is `> 0 or == -1` and not "is set". Upstream tests ported from `test_completion.py:78,281-308` and `test_completion_error.py:615-625`; the value assertions check that a position's entries are a subset of ONE `log_softmax` distribution (`sum(exp) <= 1`), which an array of zeros -- what a dropped payload looks like -- fails by summing to the entry count. Six reachability mutations, each printing `compile_rc` and `git diff --stat` and each restored byte-for-byte against a pre-taken sha256: five RED, and M6 (deleting the `ClampPromptLogprobs` CALL SITE) stayed GREEN and is owed as [#1817](https://github.com/mudler/vllm.cpp/issues/1817). TWO ARMS OF THE UPSTREAM TEST ARE DELIBERATELY NOT PORTED and the spec says why: `-1` is SERVED here rather than 400'd, because upstream compares the widened vocabulary size against a `--max-logprobs` default of 20 that our `ModelConfig` does not carry; and the `choices[1]` arm cannot run over the socket because `AsyncLLM` never fans out `n>1` ([#1816](https://github.com/mudler/vllm.cpp/issues/1816), measured, not assumed -- the same body without `prompt_logprobs` also returns one choice), so that property is gated over the SYNC `LLMEngine` instead. OpenAI `echo` stays this row's named residual and the row stays `ACTIVE`. Spec [prompt-logprobs.md](../specs/prompt-logprobs.md) `# W2` | bug | -| [#1816](https://github.com/mudler/vllm.cpp/issues/1816) | — | **`AsyncLLM` -- the engine every OpenAI HTTP route runs on -- never fans a request out into `n` children, so every `n > 1` request to the production server is silently served as `n = 1`.** `LLMEngine` does fan out (`llm_engine.cpp:151` `FanOutParallelSampling`, 1:1 `llm_engine.py:280-291`: a shared `ParentRequest`, `n` children named `{idx}_{parent}` with `n == 1` params and `seed + idx`, aggregated back into one `RequestOutput`), and all three `AsyncLLM::add_request` overloads instead pass `request_index=0, parent=nullptr` at `async_llm.cpp:79,120,278`. Measured over a real socket on `bacb71109`: `{"n":2,"temperature":1.0,"seed":7}` returns ONE choice and `completion_tokens` counts one sequence; the identical body without `prompt_logprobs` behaves the same, so the fan-out is the variable. The gap stayed invisible because the covered engine is not the served one -- `test_serving.cpp` ("serving_completion: n>1 returns n indexed, deterministic choices") gates the property over the SYNC engine. `best_of` is affected too: `serving_completion.cpp` asks the engine for `best_of` children and ranks them with `SelectBestOf`, and there is nothing to rank. Found while landing [#1815](https://github.com/mudler/vllm.cpp/issues/1815) and NOT fixed there -- porting the fan-out onto `AsyncLLM` touches the abort path and the streaming `RequestOutputKind` handling, so it needs its own row, spec and fresh review rather than an in-flow repair. Owed under [prompt-logprobs.md](../specs/prompt-logprobs.md) `## Owed` | bug | -| [#1817](https://github.com/mudler/vllm.cpp/issues/1817) | — | **The two `ClampPromptLogprobs` call sites are reached but not measured: deleting one keeps the focused gate green.** Measured while landing [#1815](https://github.com/mudler/vllm.cpp/issues/1815) as mutation M6 -- removing `ClampPromptLogprobs(prompt_logprobs);` from `serving_completion.cpp` compiles (`compile_rc=0`, deletion confirmed by `git diff --stat`) and leaves `test_openai_api_server` at 70/70 and `test_openai_serving` at 59/59. The FUNCTION is gated directly (`test_protocol.cpp`, "ClampPromptLogprobs rewrites -inf to -9999.0 in place"); the call sites are not, so a refactor could drop one silently. The cause is the fixture rather than the test: prompt logprobs come from raw prompt logits with no sampling mask applied, and a `log_softmax` over finite float32 logits does not underflow to `-inf`, so nothing in the CPU tier produces the value the clamp exists for. Repair is a seam that lets a test hand `OpenAIServingCompletion` a `RequestOutput` carrying `-inf`, or a fixture whose prompt logits contain one -- NOT a widened assertion, because a green-on-deletion gate is the defect. Owed under [prompt-logprobs.md](../specs/prompt-logprobs.md) `## Owed` | bug | -| [#1812](https://github.com/mudler/vllm.cpp/issues/1812) | `KERNEL-ACCEL-PROVIDER-DECLINE-EXACT` | **The two exact `declines == 1` routing assertions in `tests/vt/test_ops_attention_cross.cpp` are the only thing in the tree that pins the CUDA half of [#1584](https://github.com/mudler/vllm.cpp/issues/1584)'s repair, and the run CI performs cannot see it move.** Measured on `thor:gpu0` (sm_110) on 2026-08-23 while discharging [#1692](https://github.com/mudler/vllm.cpp/issues/1692), `rc` jobs `d452b91f` and `43a27be9` at `bacb71109`. Two mutations, each rebuilt and rerun on the device: reintroducing the #1584 double count in the seam (`GetOpFallbackUncounted` delegating with `/*count=*/true`), and reverting this row's whole CUDA edit (`GetOpFallbackUncounted` -> `GetOpFallback` in `src/vt/cuda/cuda_attention_cross.cu`). **Both leave the full `test_ops_attention_cross` GREEN at 20 cases / 156 assertions**, and both turn the two cases RED with `CHECK( 2 == 1 )` under a per-case `-tc=` filter. `tests/CMakeLists.txt` registers ONE ctest entry per suite, so the second row of that pair is the consequence stated plainly: the CUDA call site can be reverted to the counting resolver and every gate in this repository stays green. **Why the full run is blind was measured, not reasoned:** the reachability mutation reds `test_ops_attention_cross.cpp:326` and `:334` first -- plain `attention-cross:` cases that predate the blocked provider, assert nothing about `declines`, and run geometries `BlockedShape` rejects, so each of them resolves `BlockedFallback()`'s function-local static OUTSIDE any counted window and a later `ResetOpProviderStats` erases the extra increment. The suite warms itself, exactly as #1555's deleted `WarmDeclineOnce` warmed it by hand. The file's own comment is accurate -- exact "standalone and under `-tc=`" -- but nothing executes them that way, and `.agents/specs/op-provider-decline-exact.md` §9 said the opposite until this measurement corrected it in place. Candidate repairs: per-case ctest entries with a trailing wildcard (the depth-decoder name contains a comma and `-tc=` splits on commas, so a bare filter matches nothing and prints `0 cases ran` + `SUCCESS!`) and an asserted `test cases: 1`; or a way to reset the static between cases, which changes production code to suit a gate. Listed as O3 under `## Owed` in [op-provider-decline-exact.md](../specs/op-provider-decline-exact.md) | bug | -| [#1804](https://github.com/mudler/vllm.cpp/issues/1804) | `MODEL-MM-dots3-note-dots3-note-for-causal-lm` | **`.agents/specs/dots3-note.md` §4 item 6 read `is_neox_style=False` as belonging to the dots3-note SLIDING rope ONLY, and it belongs to both MLA ropes** — the sentence would have sent a W3 implementer to rotate the 13 full-attention layers split-half NeoX. `Dots3NoteSlidingAttention` does pass `is_neox_style=False` literally (`model.py:408` @ vLLM `origin/main` `c205726108df54bb6fbf15b19e725a4a3add2b18`), which is the half W0 read; `Dots3NoteFullAttention` (`model.py:219`) inherits the SAME hard-coded value from `deepseek_v2.py`::`DeepseekV2MLAAttention.__init__` (`:1093-1098`). So the two geometries do NOT differ on the RoPE layout at all — they differ on the THETA, `swa_rope_theta` 5e4 on 33 layers against `rope_theta` 8e7 on 13. The polarity that DOES flip is the INDEXER's, and that is §4 trap 2's point rather than item 6's: `deepseek_v2.py:1148` sets the indexer rope to `is_neox_style = not indexer_rope_interleave`, so at DeepSeek-V3.2's absent-key default the indexer runs NeoX beside an MLA rope that is GPT-J, and `indexer_rope_interleave = True` (`configs/dots3_note.py:23`) is what makes dots3-note's two agree. Being wrong in this direction is SILENT: the same 64 coordinates are rotated either way, so nothing changes shape and nothing throws — the §4 defect class exactly, on a row that spec §6.4 says has no oracle anywhere to catch it. FIXED IN FLOW on `row/MODEL-MM-dots3-note-W1`: §4 item 6 corrected in place and it says what it used to say (`main` is never rewritten), `ParseDots3NoteParams` resolves `rope_is_neox_style = false` on BOTH geometries with the two citations beside it, and `tests/vllm/models/test_dots3_note_scaffold.cpp` asserts both plus the indexer's agreement with them. The assertion was captured RED against the NeoX reading, on an arm that compiled and ran, before the corrected value existed | bug | -| [#1802](https://github.com/mudler/vllm.cpp/issues/1802) | `BACKEND-CUDA-SM110` | **Re-measuring the sm_110 `ctest` baseline at `6756f9131` (2026-08-23, job `8bf39567-9334-4f7e-aa27-43a2aa867bb7`, artifacts `/mnt/nas_share/rc/thor-w05-955/out/`) found six names red that were green at `0764ded2b`, none with an owner**, plus two standing entries whose recorded cause no longer holds. Arrivals, all mode `Failed`: `test_cuda_ops` (`:106` `CHECK(bad == 0)` -> `6 == 0` and `7 == 0`, silu_and_mul vs CPU; also red on GB10 at 439/440 against 438/440 here, [#907](https://github.com/mudler/vllm.cpp/issues/907)), `test_backend_cross_device` (`:2063` MoeSiluMul and `:2601` MoE combine/gate vs the CPU oracle, 80205/80207), `test_llama_embedding_fold` (`:254` engine vs direct at `epsilon(1e-5)`), `test_mtp_depth` (`:738` `CHECK(st.capture_shapes == 0)`), `test_qwen3_dflash2_draft` (`:2574` `CHECK(r.generate_threw.empty())`), and `test_ops_attention_dense_fa2` (`:692` `CHECK(Mismatches(on, ref) > 0)` -> `0 > 0`, which reads as a TEST arch-assumption: the case asserts the knob-ON path DIFFERS from the scalar reference, and with `fa2` DISABLED for `[110]` they are the same kernel). Four of the six are CUDA-vs-CPU numeric comparisons off by about one ulp, so they MAY be one sm_110 difference wearing four names -- a hypothesis from the logs, not a measurement. Also folded in: `test_gguf_device_fit_reach`, red and unattributed since 2026-08-15 (`:463`, moved from `:278`, assertion unchanged); and `test_serve_low_tools`, whose `(name, mode)` pair never moved while its CAUSE changed completely -- `73ada0df8` (#1661/#1662) fixed the [#961](https://github.com/mudler/vllm.cpp/issues/961) guard so it skips, the string `shellcheck` appears nowhere in the run's log, and the failure is now four `tests/tools/test_dflash2_speed_harness.py` `ShellDriverTest` cases of 517. The CONTROL was rerun and now falsifies the old conclusion: the baseline reads `FAILED (failures=3, errors=1, skipped=1)` and the control with `shellcheck` 0.9.0 installed reads `FAILED (failures=3, errors=1)` -- the same four cases, differing only in the vanished skip. **That vanished `skipped=1` is a second proof**: the one test that skipped in the baseline is the `shellcheck` guard itself, so with the binary present it stopped skipping and PASSED, showing the instrument was the only thing the install changed. #961 was CLOSED COMPLETED on 2026-08-23 acting on that prompt -- `73ada0df8` fixed its guard while referencing the sibling filing #1661/#1662, which had left #961 orphaned rather than resolved. Not in scope here, because they have owners and did not worsen: the four FA-2 refusals, `test_platform`/`test_op_parity` hardcoding GB10, `test_linear_method` and the three `qwen3_5_gdn_spec_routing` tests (#907), the FP8 pair ([#1725](https://github.com/mudler/vllm.cpp/issues/1725), which improved `SEGFAULT` -> `Failed`), and [#962](https://github.com/mudler/vllm.cpp/issues/962), which reproduced byte-identically at `bitdiff=15/32768`. Table and diff in [`environment.md`](../environment.md); the re-measure debt itself is [#955](https://github.com/mudler/vllm.cpp/issues/955) | bug | -| [#1829](https://github.com/mudler/vllm.cpp/issues/1829) | `MODEL-MM-dots3-note` | **`src/vllm/model_executor/models/dots3_note.h:282` declares `[[noreturn]] static ForwardLogits ForwardDevice(...)` -- `[[noreturn]]` on a non-void return type -- and MSVC's C4646 plus warnings-as-errors turns that into `error C2220` at `dots3_note.cpp(606,31)`, so the whole `vllm` project fails to COMPILE on Windows.** GCC and Clang accept the declaration silently, so `build-test-cpu`, `build-test-cpu-arm64`, `build-newest-gcc` and `verify (cpu)` are all green on the same commit. Found on the `windows-msvc-cpu` job of #1821, whose own diff is four files and none of them this one; the declaration arrived with `849a7dd73` (#1805) which is an ancestor of `af320abb2`, so it is INHERITED and every pull request branched from `main` since carries it. It landed unseen because `windows-msvc-cpu` and `windows-msvc-vulkan` are PULL-REQUEST-ONLY jobs with no `main` baseline, which is the second half of the defect: `main` receives no Windows verdict at all. **NOT #584** -- that is the runtime `exit -1073740791` STATUS_STACK_BUFFER_OVERRUN in `test_openai_api_server.exe`, and a reader who stops at the job name will wave this real break through as the known one. Not repaired in flow by #1821: `ForwardDevice` overrides into a registry hook and cannot simply become `void`, so the fix is a semantic decision in an actively-developed file the `MODEL-MM-dots3-note` row owns, and #1821 has no MSVC to verify one against | bug | -| [#1809](https://github.com/mudler/vllm.cpp/issues/1809) | `GATE-CI-ENFORCEMENT-FLOOR` | **`commit-protocol-tag` and `documentation-checkpoint` have been red on `main` since 2026-08-13 and cannot recover on their own: their walk bases on the head of the last SUCCESSFUL push run, so one unrepairable commit freezes the base and every later push re-walks the same violations over a range one commit wider.** Re-measured 2026-08-23 at `bacb71109`: `LAST_GREEN` is `fafa16f0f` of 2026-08-13T01:53, the range is **499** first-parent commits with **0** merges, and **41 distinct commits** in it can never be repaired because they are on `main` and adding a trailer or a task-branch arrival needs a rewrite `AGENTS.md` forbids — **35** fail `check-commit-trailers.py --range` (of which **20** carry no `FOLLOWING_AGENTS_PROTOCOL` marker at all and the other **15** carry it 1 to 9 times in a form the strict contract rejects, so the two counts describe the same population under different contracts, not different walks) and **6** disjoint ones fail `check-role-discipline.py` for a repository change that reached `main` with no task branch; `check-now-current.py` passes over the whole range and forgives nothing. The `LAST_GREEN` base is RIGHT about what it was built for and is not reverted: `github.event.before` skips a cancelled run's commits and nothing re-covers them, which is what lets the push lane be latest-only ([#822](https://github.com/mudler/vllm.cpp/issues/822), [#863](https://github.com/mudler/vllm.cpp/issues/863)), and cancelled runs are common here ([#1285](https://github.com/mudler/vllm.cpp/issues/1285)). FIXED IN FLOW with a recorded **enforcement floor** in `scripts/ci-enforcement-floor.txt` — one commit the walk never goes behind — taken as the newer of `LAST_GREEN` and the floor by ANCESTRY (`git merge-base --is-ancestor`, never by commit date, which is author-controlled and can go backwards). Losslessness survives because the floor is a lower clamp on an otherwise unchanged base: while the floor sits behind `LAST_GREEN`, which is the steady state, the resolved base is byte-identical to today's, and the only window that skips anything is `LAST_GREEN..floor` right after an advance — bounded, reviewed, and enumerated in the spec. The four byte-similar copies of the selection in `.github/workflows/ci.yml` collapse into `scripts/ci-walk-base.py`, which is the point of the extraction: of those four inline blocks exactly ONE was ever replayed by a test (`test_main_baseline.py::AgentRecordDiffRangeTests`, on `agent-record`'s body, under a shim that stubs every `python3` call and therefore pins which checker ran with which range string rather than the base rule), and the other three were executed by nothing, so the property most likely to break silently now has a test that replays a cancelled-run sequence against a real throwaway repository with the naive `before` base as its positive control. A per-commit exemption list was REJECTED — `AGENTS.md` has no waiver registry, the mechanism would have to be built three more times in checkers that have no exemption concept, and it never shrinks the walk, so the range grows by one per merge forever. This NARROWS enforcement by 41 commits and says so: enforcing on an immutable landed commit is a permanent red rather than a gate, because no contributor action can clear it, and a permanent red is read by nobody ([#1722](https://github.com/mudler/vllm.cpp/issues/1722)'s effect). No assertion is deleted and no checker is modified; only the base of the walk moves, and a new violation after the floor still reds, proved by a scratch-commit mutation rather than by reading the diff. Spec [ci-enforcement-floor.md](../specs/ci-enforcement-floor.md) | bug | -| [#1823](https://github.com/mudler/vllm.cpp/issues/1823) | `BACKEND-ATTN-REGISTRY` | **`FlashAttentionBackend::supports_compute_capability` is upstream's NVIDIA `capability >= (8,0)` (`include/vllm/v1/attention/backend.h:401-404`, `flash_attn.py:200-202`), and `MetalPlatform::get_device_capability` (`src/vllm/platforms/metal.cpp:35-38`) answers with the MTLGPUFamilyApple GENERATION — deliberately, with a comment saying `has_device_capability(N, 0)` means "Apple family >= N". `validate_configuration` at `src/vllm/v1/attention/backend.cpp:197` compares the two whenever `capability.present()`, so an SM-version bar is applied to an Apple family number.** Apple family 9 on the M4 gate box clears it by coincidence; a GitHub `macos-15` runner reports lower, FLASH_ATTN — the ONLY entry in `MetalPlatform::get_attn_backend_priority()` — is refused, and `SelectAttentionBackendName` throws. MEASURED 2026-08-23 on the first execution of `test_metal_backend` since `7020de936` (v0.0.2, 2026-08-11), run [32668677681](https://github.com/mudler/vllm.cpp/actions/runs/32668677681) at `744abeca3`: `tests/vt/test_metal_backend.cpp:170` ERROR `"No valid attention backend for device type 2 from {FLASH_ATTN: [compute capability not supported]}"`, suite 26 cases / 25 passed / **1 failed** / 3 skipped, 112336 assertions, `Status: FAILURE!`. **Red since `369ea7fd4` (2026-08-19), which is NOT an ancestor of `7020de936`** — four days in the exposure window [#1765](https://github.com/mudler/vllm.cpp/issues/1765) measured, with no lane able to run the case. **The prose above `backend.cpp:197` states the premise that is false**: it argues `DeviceCapability::present()` is already false for every platform that cannot answer, and Metal answers in a different unit; `vulkan.cpp` and `tenstorrent.cpp` need the same check before any repair. **Not fixed in flow**: every candidate repair (drop the predicate for non-SM platforms, stop reporting `present()` on Metal — which `test_metal_backend.cpp:153-154` asserts on purpose — or make the backend device-type aware) changes what selection means for kCPU/kMETAL/kVULKAN/kTENSTORRENT, so it takes the normal row, spec and fresh-review path. Until then `macos-metal-mlx` is RED on `main`, deliberately not hidden behind `continue-on-error` or a skip. Listed under `## Found in flow, filed, not fixed here` in [attn-validate-configuration.md](../specs/attn-validate-configuration.md) | bug | -| [#1819](https://github.com/mudler/vllm.cpp/issues/1819) | `ENG-MM-INPUT-PIPELINE` | **`scripts/mm/tower_skip_rss.sh` configured both build directories with `-DVLLM_CPP_BUILD_EXAMPLES=OFF` and then ran `ninja -C "$d" -j 4 vllm-server`, and `vllm-server` is an `examples/` target** -- the `OUTPUT_NAME` of `server` (`examples/CMakeLists.txt:91,108`), in a directory the root `CMakeLists.txt:2828` adds only under `if(VLLM_CPP_BUILD_EXAMPLES)`. Reproduced with the harness's own flags: the configure returns 0 and `ninja` answers `unknown target 'vllm-server'`, so the run `exit 4`s at arm A before any RSS exists. The block landed on `main` in `bacb71109` (#1364) and had never been executed. **Nothing could catch it**: `tests/scripts/test_tower_skip_rss_report.py` covers `--report-only`, `--check-source` and `--stage-check` -- every path needing no checkpoint -- while the configure, the build, `run_arm`, the `/health` poll and the kill/wait only ever run on a leased box, so the suite was 41/41 green over a harness that could not build its own binary. FIXED IN FLOW: `-DVLLM_CPP_BUILD_EXAMPLES=ON` (measured: `ninja -j 4 vllm-server` then returns 0 and writes `/examples/vllm-server`, the one file of that name in the tree and the path `docs/USAGE.md:54,95,128,204` names), the binary is NAMED rather than found by a `find` piped into `head -1`, and a new `--dry-run` prints the `cmake`/`ninja`/`run_arm` invocations out of the same variables the run issues them from and asserts that CMake defines the requested target under those flags -- statically, so it builds nothing and runs in CI, plus a live `ninja -t targets` prong on an already-configured tree that skips BY NAME when there is none. `DryRunTests` runs it against the script as committed, which is the case that reds on this defect, and against scratch copies with the flag flipped OFF, the flag dropped, and the binary path pointed away from where CMake writes it | bug | -| [#1647](https://github.com/mudler/vllm.cpp/issues/1647) | `LOAD-MODELOPT-NVFP4-BORROW` | **`LoadNvfp4AnyNaming`'s ModelOpt branch never attempted the ENG-LOAD-DIRECT-UPLOAD borrow, so `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121`'s 21 GiB of packed weights landed as ANONYMOUS heap and then got a second, device copy on a box where the two come out of the same 119 GiB.** Measured on the committed tree at `27d8bfa70`, 2026-08-23, on `dgx:gpu0` (GB10, 119 GiB UNIFIED): `VT_LOAD_STATS=1` reads `host_copy=16.394 GiB borrowed=2.373 GiB device_upload=0.000 GiB`, which FALSIFIES the reflex diagnosis — the weights reach the host PACKED, not dequantized to bf16, and nothing here chases a bf16 materialization. What it shows instead is 16.394 GiB of anonymous heap where upstream holds reclaimable PAGE CACHE. A small run (`--kv-cache-memory 512MiB --max-num-seqs 1 --max-model-len 1024`) loads with 85 GiB left; the serving shape (`8GiB` KV, `--max-num-seqs 4`, `--max-model-len 4096`, `fp8` KV) collapses `avail 37 GiB -> 28 GiB in ONE second` and, unguarded, reaches 0 and reboots the host with NVRM `NV_ERR_NO_MEMORY ... _memdescAllocInternal` — five times. vLLM on the SAME checkpoint and the SAME box at `--gpu-memory-utilization 0.70` sits at 71-83 GiB available and serves. MECHANISM, derived from the tree before the fix was written: `LoadCtNvfp4Raw` has carried `BorrowStTensorBytes`-then-`MakeOwned` since ENG-LOAD-DIRECT-UPLOAD, and the ModelOpt sibling goes straight to `MakeOwned` + `std::memcpy` for BOTH `packed` and `scale`; this checkpoint declares `quant_method: "modelopt"`, so all 193 NVFP4 modules take that branch. The expensive consequence is the second one: `AdoptDeviceBytesAsHost` is gated on `w.mmap_src != nullptr && w.bytes.borrowed()`, so `ResidentNvfp4`'s post-upload adoption — which is what releases the source pages and, where device memory is host-addressable, collapses the host and device copies onto one buffer — was a SILENT NO-OP for every ModelOpt weight. `ResidentNvfp4` has no aliasing branch (unlike `ResidentWeight`), so the device copy is a genuine second allocation: ~42 GiB of a 119 GiB pool before one KV byte is reserved. FIXED IN FLOW: the ModelOpt branch gets the identical borrow-then-fallback shape, byte-size `VT_CHECK`s and fallback preserved — a widening of an already-gated path to a second checkpoint SPELLING, not a new policy. UPSTREAM AT THE PIN `5559679229bc961848b121ccdeaa8fa5d79bec98`, read in a local checkout at that exact SHA: vLLM allocates the FINAL PACKED parameter ON THE DEVICE before opening any file (`base_loader.py:52-58` `with target_device:` around `initialize_model`, `modelopt.py:1313-1345` uint8 packed + `float8_e4m3fn` scales with NO `device=`, `:483-491` for the FP8 half) and streams the shard into it one tensor at a time through an mmap generator (`weight_utils.py:969-974`, consumed at `:1247` `param.data.copy_`), so its host-side residency is one loop iteration of page cache. A SECOND EDIT reconciles two predicates that disagree about one device: `DirectDeviceLoadEligible` required `!platform.is_unified_memory()`, which is true on GB10, while `needs_weight_staging()` is unconditionally true on CUDA and `src/vllm/platforms/cuda.cpp:85` says why in its own comment — "regardless of GB10 being physically unified" — and `include/vllm/platforms/interface.h:294-298` records that `is_unified_memory()` "answers the OPPOSITE question". `ResidentNvfp4` proves the staging predicate right by construction. THE SECOND EDIT DOES NOT FIX THE OOM and the spec says so rather than implying otherwise: `StageAndReleaseLoadedDense` is reached only while `IsPlainBf16Qwen3_5Dense` holds, which this checkpoint fails at its first NVFP4 projection, and `ReleaseResidentQwen3_5DenseHostWeights` walks only `OwnedTensor`s. Its measurable effect is a plain bf16 dense checkpoint on unified memory. It is not the use-after-free the release site's comment warns about: that site asks the invariant `HostMirrorIsRedundant` (`d_dev != nullptr`), and the aliasing arm leaves `d_dev` null. NOT CLOSED: #1647's own `## Owed` — the DFlash2 startup instrumentation, the double dequant-B question, a memory bound that refuses before the global OOM killer takes `coredns` with it, and the re-run — is live and listed under `## Owed` in the spec, together with upstream's integrated-GPU memory accounting (`mem_utils.py:148-155`, `:54-83`), which this tree has nowhere, and the allocate-device-first polarity that is the faithful long-term target. Every gate here is CPU-only: this session had no GPU and no fleet access, so the `avail` curve at this head is owed. Spec [`load-modelopt-nvfp4-borrow.md`](../specs/load-modelopt-nvfp4-borrow.md) | bug | -| [#1791](https://github.com/mudler/vllm.cpp/issues/1791) | `SPEC-DSPARK` | **`scripts/dspark-paired-e2e.sh`'s `settle()` can never break early, so a wait for the GPU to drain always spends its full 360 s however fast the box actually drains.** The same idiom as [#1734](https://github.com/mudler/vllm.cpp/issues/1734), found by sweeping `scripts/` for it: `grep -c .` with an `\|\| echo 0` fallback makes `$n` the two-line string `0\n0`, so `[ "$n" -eq 0 ] && break` answers `integer expression expected` and returns 2 instead of deciding. The BUSY half of the guard works -- a positive count exits 0 and the fallback does not fire -- so only the FREE half is dead, and the failure is in the safe direction, which is why it was paid in silence. This is the FIFTH diagnosis of the idiom in this tree: `scripts/cpu-x86-llamacpp-floor.sh` already carries the removal and the reason in a comment, and a comment in one file is not reachable from another. FIXED IN FLOW with #1734: `\|\| true` keeps grep's own `0` and swallows only its status. The recurrence gate is `TheIdiomIsGoneFromEveryShellScript` in `tests/scripts/test_ltx25_ab_memwatch.py`, which sweeps every `scripts/*.sh` for a counting `grep`/`pgrep` paired with an `\|\| echo` fallback outside a comment; run against `27d8bfa70` it names all three live instances, this one included. It is a TRIPWIRE and says so: it reads text, and a `wc -l` with the same fallback walks past it | bug | -| [#1796](https://github.com/mudler/vllm.cpp/issues/1796) | `SPEC-DFLASH2` | **[#1456](https://github.com/mudler/vllm.cpp/issues/1456)'s conclusion is retracted AT THE ARTIFACT, off-GPU, and the DFlash2 speed ratio's denominator inherits the retraction: `0.8016987337853048` MAY BE FLATTERING US.** #1456 measured a real source build at `CUDA_ARCHS=12.0` emitting `sm_80` for `_vllm_fa2_C` and `sm_75` for `_vllm_fa3_C`, and concluded the GB10 oracle has no `FLASH_ATTN` denominator. **The arch measurement stands; the conclusion does not, and no lease was needed to show it.** Both staged oracle wheels were opened with `zipfile` and walked as fatbinaries on the CPU dev box: `0.1.dev1+g66e5414c6` (this row's oracle, `/mnt/nas_share/rc/oracle-dflash2/`) and `0.1.dev1+g555967922` (the parity pin, `/mnt/nas_share/rc/oracle-vllm/`) each carry 76 FA2 fatbinaries and 192 FA3 fatbinaries, and **every one holds a PTX image beside its SASS image** -- arch 80 for FA2, arch 75 for FA3. The first FA2 PTX payload is zstd and decompresses to `.version 9.0` / `.target sm_80` for `flash_fwd_hdim128_bf16_causal_sm80`. That is the `+PTX` half of `FA2_ARCHS "8.0+PTX"` and it is the mechanism by which the module CAN reach sm_121 -- the artifact establishes a necessary condition, forward-JITtable code being shipped, and that the JIT then ran is an inference from that plus a run selecting `FLASH_ATTN` and generating; `cudaErrorUnsupportedPtxVersion` is the OPPOSITE failure, PTX ISA newer than the driver, which `.version 9.0` under driver 580.173.02 is not. vLLM says it directly too: `FlashAttentionBackend.supports_compute_capability` returns `capability >= DeviceCapability(8, 0)` (`vllm/v1/attention/backends/flash_attn.py:251-252` in the staged wheel). **THE DOWNSTREAM IS THE THING TO CHASE.** #1456's body records that the DFlash2 speed gate's denominator "will be vLLM pinned to `TRITON_ATTN`, by developer decision on 2026-08-20", explicitly because of that conclusion. AGENTS.md requires vLLM's PRODUCTION configuration as the denominator, and on this box vLLM's own auto-selection picks `FLASH_ATTN`: the 0.8017 gate run's log carries BOTH paths in one process, the forced branch taking `TRITON_ATTN` for the 27B target (`cuda.py:426`) and the auto branch taking `FLASH_ATTN` out of four valid backends (`cuda.py:486`), in `/mnt/nas_share/rc/dflash2-1673/out-n1673b/m-gate.log`. **So if `TRITON_ATTN` is the slower backend -- which vLLM's own priority ordering IMPLIES rather than states, and which nothing here measures -- then 16.279 tok/s is too LOW and the ratio is too HIGH. An error in our own favour is the one nobody chases, so it is recorded with its sign.** The ratio is NOT withdrawn and no denominator is substituted, because a wave must not replace a denominator the developer declared. **The five FA layers of [#1685](https://github.com/mudler/vllm.cpp/issues/1685) fall out of the same read**: `vllm/v1/worker/gpu/spec_decode/dflash/utils.py:31-46` sets the draft's `backend=speculative_config.attention_backend` UNCONDITIONALLY, so a harness that sets only the engine backend leaves the draft at `None` and the draft re-runs auto-selection; the siblings `dspark/utils.py:24-28` and `gemma4/speculator.py:66-89` both carry the target's backend through, and dspark's comment names this hazard by name. NOT reconciled in flow, and the reason is precise: what remains is ONE lease running vLLM against itself on this workload with `attention_backend=FLASH_ATTN` against `TRITON_ATTN`, each read back off the built engine, and **no GPU lease is authorised for this session**, so it is PENDING. Recommended collapse, recommended and not taken: #1456, [#1679](https://github.com/mudler/vllm.cpp/issues/1679) and #1685 observe ONE mechanism on two wheels, and #1685 is the one attached to a live number and should survive. Owed under `## Owed` O33 of [the DFlash2 spec](../specs/dflash2-spec-decode.md), with caveat 5 beside the ratio in the same file | verification | -| [#1794](https://github.com/mudler/vllm.cpp/issues/1794) | `LTX25-DIT-ATTN-ARM-PARSE` | **#1551 renamed the values of `VLLM_LTX2_DIT_FLASH_ATTN` and the two #1549-era LTX-2.5 A/B harnesses kept exporting the old ones, so three committed arms name a rung they do not select -- and one of the three says nothing.** At [#1549](https://github.com/mudler/vllm.cpp/issues/1549) the knob was BINARY: `=0` selected `vt::Attention` and every other value, unset included, selected `vt::AttentionDenseFlash`. [#1551](https://github.com/mudler/vllm.cpp/issues/1551) made it THREE-WAY, moved the unset default up a rung to `vt::AttentionDenseFa2` and gave the flash rung the exact spelling `flash`; it touched neither harness. Read at `27d8bfa70`: `scripts/ltx25-dit-attn-flash-pixel-ab.sh` exports `=1` for both its `flash` and its `flash-ctl` arms, and `scripts/ltx25-dit-attn-flash-ab.sh` leaves the variable UNSET for its `flash` arm -- all three of which selected `AttentionDenseFlash` before #1551 and select `AttentionDenseFa2` after it. `scripts/ltx25-dit-attn-fa2-hd128-ab.sh` is CORRECT and is the contrast that makes the other two legible: written after the rename, it exports `flash` / `""` / `0` and carries `assert_arm_op`. THE ASYMMETRY IS THE ISSUE: the pixel harness's `arm_report` counts `op=18` and `op=21` in the arm's own log and exits 46 on `ROUTING_BAD`, so its two `=1` arms abort -- after the render, an hour of a four-hour lease spent to learn a literal went stale -- while `ltx25-dit-attn-flash-ab.sh` has NO such assertion at all: its phase `[F]` PRINTS the op-provider selections and asserts nothing, so its `flash` arm renders FA-2, is reduced under the label `flash`, and the ratio it publishes is FA-2 against naive while the file's own header says flash against naive. THE RECORDED NUMBERS ARE NOT INVALIDATED: `.agents/specs/ltx25-dit-attn-flash.md` section 10 records `knob=1` announcing `op=21` and `ROUTING_OK=flash`, so the binary that ran it is one where `=1` still selected flash; what is wrong is the file as it stands against the binary `main` builds today. WHY A GATE AND NOT A REVIEWER: both harnesses were TRUE when written and both became false in a commit that touched neither file, because nothing connected the literal in a shell script to the literal in the C++ dispatch. FIXED IN FLOW with the red-first case that would have caught it -- `tests/scripts/test_ltx2_dit_attn_knob_arms.py` reads the accepted set out of `ltx2_device.cpp`'s own `std::strcmp(arm, "...")` calls rather than restating it, reads each harness's arm invocations, and asserts BOTH that every value is one the dispatch parses AND that every arm selects the rung its LABEL claims; the second half is the one that matters, because `unset` is a perfectly valid value and was still the wrong arm for a whole row. Registered on the preflight and CI record lanes, and it asserts both registrations itself. Found while fixing [#1751](https://github.com/mudler/vllm.cpp/issues/1751), which is the same defect one level down. Spec [`ltx25-dit-attn-arm-parse.md`](../specs/ltx25-dit-attn-arm-parse.md) | bug | -| [#1832](https://github.com/mudler/vllm.cpp/issues/1832) | `SGLANG-ORACLE-LEASE-WHEEL` | **The SGLang manifest's file count `3338` is gated only against itself, and appears in no executing code.** Raised by the fresh review of PR #1831 and NOT repaired there: the number and the claim are W1's (`727efb39c`), and re-deriving them needs an `rc` job on `dgx:gpu0`, which is outside W2's scope. `3338` is quoted as MEASURED in three records -- `.agents/environment.md`, `.agents/oracles/sglang.md`, and the `SGLANG-ORACLE-LEASE-WHEEL` row of `.agents/sglang-matrix.md` ("**3338 of 3338** manifest files, 0 missing, 0 extra, 0 differing") -- while `grep -rn '3338' scripts/ tests/scripts/ .github/` returns `rc=1`. The only test that touches the count is `tests/scripts/test_sglang_lease_identity.py:54-55`, which asserts `manifest["file_count"] == len(manifest["files"])`: a self-consistency check of one JSON document, which cannot see whether 3338 is the count of the real installed tree. MUTATION, run at `85c247580`: drop `sglang/README.md` from `files` and decrement `file_count` to 3337 -- the exact shape of a mis-generated manifest -- and the suite reads `Ran 14 tests ... OK`, `rc=0`. Tree restored byte-for-byte. This matters because `sglang.__commit_id__` is `None` in the published wheel, so the manifest is the ONLY identity assertion available for this oracle, and a wrong manifest makes `IDENTITY_RC=0` a tautology one level up. A checker reading the committed JSON cannot repair it; the repair is a re-derivation -- a second independent install that REGENERATES the manifest and diffs it against the committed one. Until that runs the honest record is "3338 files, from one generation run on 2026-08-19, not independently re-derived". Listed under `## Owed` in [sglang-wheel-in-lease.md](../specs/sglang-wheel-in-lease.md) | bug | -| [#1833](https://github.com/mudler/vllm.cpp/issues/1833) | `SGLANG-ORACLE-LEASE-WHEEL` | **Both registrations of `test_sglang_lease_identity` are deletable at `rc=0`, and the control shows the `SUITES` array itself is ungated.** Raised by the fresh review of PR #1831 and NOT repaired there: the control proves it is a property of the array rather than anything W2 introduced, and repairing `check-test-registration.py` is a semantic checker change owing its own row, spec and red-first evidence. The suite is registered TWICE, deliberately: `scripts/agent-preflight.sh:176` and `.github/workflows/ci.yml:341`. Two mutations at `85c247580`, each proven applied (non-empty diffstat, mutated file still parses) and each restored byte-for-byte: deleting the `SUITES` entry leaves `python3 scripts/check-test-registration.py` at `rc=0`, and deleting the whole 11-line CI step leaves it at `rc=0` -- `grep -c` on the workflow returns 0 and `yaml.safe_load` still parses it. CONTROL: deleting the unrelated `test_tower_skip_rss_report` from the same array behaves identically, `rc=0`. So ANY entry can be removed from `SUITES` with no gate noticing, and the "Registered in TWO places, deliberately" pattern that `.github/workflows/ci.yml:339-343` documents buys no protection in either direction. NOT covered by the neighbours: [#408](https://github.com/mudler/vllm.cpp/issues/408) names suites executed by NOTHING and correctly diagnoses `REQUIRED_TESTS` as a self-guard rather than a population guard, and [#1730](https://github.com/mudler/vllm.cpp/issues/1730) names a suite with ONE registration absent from `SUITES`; neither states that deleting a `SUITES` entry is itself ungated, which is what the control measures. Triage may prefer to fold this into #408. Repair, when taken: give the checker a POPULATION rule -- every `tests/scripts/test_*.py` must appear in `SUITES` and in a CI lane, red when either falls away -- which will red on the twelve suites #408 lists, and that is the point. Listed under `## Owed` in [sglang-wheel-in-lease.md](../specs/sglang-wheel-in-lease.md) | bug | -| [#1824](https://github.com/mudler/vllm.cpp/issues/1824) | `SPEC-DFLASH2` | **A configured speculator forces synchronous scheduling; upstream keeps async ON for the Eagle-type family (dflash included).** `model_loader.cpp` resolves `async_scheduling_enabled_` to false whenever `resolved_spec_config_` is present, a deferral recorded at SPEC-MTP I5d ([mtp-spec-decode.md](../specs/mtp-spec-decode.md) §2.8's not-ported list). Upstream's polarity at the pin (`vllm/config/vllm.py:1064-1112`) disables async only for a method OUTSIDE `EagleModelTypes ∪ NgramGPUTypes ∪ {"dspark"}` — and `"dflash"` (which DFlash2 rides) and every MTP type are Eagle-type. At c1 spec decode (~360 steps / 2048 tokens) every host-side scheduling cost is serialized into each step, the largest named host-side divergence in the [#1574](https://github.com/mudler/vllm.cpp/issues/1574) gap. W7 under `SPEC-DFLASH2` ports the draft-in-output flow (AsyncScheduler `-1` placeholders, worker-side fill, `update_draft_token_ids_in_output`, the `async_tokens_to_discard` rollback guard) and flips the enable to upstream's method predicate; the GPU TPOT A/B stays owed to the operator. Spec [`spec-decode-async-scheduling.md`](../specs/spec-decode-async-scheduling.md) | feature | -| [#1754](https://github.com/mudler/vllm.cpp/issues/1754) | `GATE-CI-SITE-HUGO-LANE` | **`agent-record` is RED on `main` and on every pull request cut from it: `test_rendered_benchmark_index_links_resolve_to_emitted_pages` shells out to `hugo` and no job in `ci.yml` installs it, so `subprocess.run` raises `FileNotFoundError` before any assertion runs and `unittest` reports an ERROR.** Introduced by [#1714](https://github.com/mudler/vllm.cpp/pull/1714) at `1db7e59cf`; the scheduled baseline at `deb53c6a3` was green on this job and `08c81a892` and every scheduled run after it is red. FIXED IN FLOW by installing the renderer on the lane -- the same `peaceiris/actions-hugo@v3` action, the same `extended: true`, and the same `0.146.3` pin the `gh-pages` job publishes on -- plus `hugo version` as the step's first line and `tests/scripts/test_ci_site_lane.py` holding the two pins equal and the setup step present and ordered before the render. A SKIP GUARD WAS REJECTED AS THE FIX and is deliberately not in this change: it retires the red by arranging for the assertion to run nowhere, which [PR #1726](https://github.com/mudler/vllm.cpp/pull/1726) demonstrates -- its `agent-record` is SUCCESS having rendered no site. The guard is still wanted for a developer box without Hugo and is separately owned by [#1722](https://github.com/mudler/vllm.cpp/issues/1722) / #1726, which edits a disjoint file, so the two compose. [#1764](https://github.com/mudler/vllm.cpp/issues/1764) names this red as one of its three and is not closed by this change | bug | -| [#1828](https://github.com/mudler/vllm.cpp/issues/1828) | `GATE-CI-SITE-HUGO-LANE` | **`tests/scripts/test_check_site.py` asserts a literal `10` rendered benchmark detail links, which is a stored count of `docs/benchmarks/*.md` inside another file, and #1754's lane fix is what ARMS it**: until Hugo was installed the case never reached that line in CI. Measured on `d60692c89` with Hugo 0.146.3+extended: 10 slugs, 10 table hrefs, no duplicated target and no unlinked slug, so the relationship is a bijection and the literal is derivable at read time -- with its own non-vacuity floor, since `sorted(x) == sorted(y)` is satisfied by two empty sets. NOT FIXED IN FLOW: deriving it changes what the case asserts rather than how it spells a number, which `AGENTS.md` routes through the normal row, spec and fresh-review path, and the file already has an open pull request against the same case. Listed under `## Owed` in [gate-ci-site-hugo-lane.md](../specs/gate-ci-site-hugo-lane.md) | bug | -| [#1732](https://github.com/mudler/vllm.cpp/issues/1732) | `FIX-CUBLASLT-CAPTURE-1732` | **`cublasLtMatmulAlgoGetHeuristic` returns status 14 while a created stream is in CUDA graph capture on CUDA 13.3, so the default (graphs-on) decode path dies on its first capture step on sm_80.** Reproduced on a CMP 170HX (GA100, driver 610.57.04, CUDA 13.3 V13.3.73) at main `08c81a892`: the bf16/f32 GEMM paths in `src/vt/cuda/cuda_matmul.cu` query the heuristic per call, cuBLASLt's own trace names `Could not obtain green context information`, and a four-line standalone probe isolates the query-under-capture failure with no vllm.cpp code. `VLLM_CPP_CUDAGRAPH=0` runs the identical workload correctly (24.3 tok/s warm, deterministic). Fix: per-key heuristic cache, default on, populated by the eager warm step every slot runs before capture. Spec: `.agents/specs/fix-cublaslt-capture-heuristic.md`. | bug | -| [#1837](https://github.com/mudler/vllm.cpp/issues/1837) | `SPEC-DFLASH2` | **DFlash2 propose round-trips ~25 MB and ~10 full-queue syncs per step because the `final_out` host contract takes every step off the D13 paged+graph draft forward.** Measured 2026-08-24 on the #1574 workload: step 141.5 ms vs vLLM ~127 / SGLang ~121 at equal acceptance, and the whole remaining gap is propose overhead the in-code debt note in `runner.cpp::propose_drafts_block` already names. W8 mirrors the merged upstream `_generate_draft` (`b389ac2946`): the block forward hands out DEVICE logits+hidden, the selector runs TopK/edges/walk device-to-device, and only the K draft token ids come back. Wave spec [dflash2-device-propose.md](../specs/dflash2-device-propose.md); the GPU TPOT number is owed there, operator-run | perf | -| [#1838](https://github.com/mudler/vllm.cpp/issues/1838) | `SPEC-DFLASH2` | **The DFlash propose pre-phase ping-pongs the `[T, H*taps]` aux tap through the host (D2H + scalar cast + H2D + D2H) with three hard syncs per step, before the `t_fwd0` timer, so `VT_SPEC_TRACE` cannot even see it.** `exec_state_.spec_aux` is already a device tensor and upstream's `propose` consumes it device-side with no round trip. Fixed in the same W8 flow as #1837 (one reshape of one propose path): `CombineAuxFeaturesDevice` GEMMs the bf16 tap directly, the accepted-prefix gather is a device `IndexSelect` over indices the rejection output already determines, and the context-KV append is fed device-side; the `[spec-phase]` trace gains a `pre=` term so the phase stays attributed. Serves DFlash1 and DFlash2 alike (the pre-phase is the shared seam). Wave spec [dflash2-device-propose.md](../specs/dflash2-device-propose.md) | perf | -| [#1839](https://github.com/mudler/vllm.cpp/issues/1839) | `MUSIC3-DEPTH-DEVICE` | **The engine's call to `Music3SelectDepthArm` (`minimax_music3_speech.cpp:638`) is reachable but not gated, and [#1131](https://github.com/mudler/vllm.cpp/issues/1131) no longer covers it: #1131 named both device-arm twins and row `MUSIC3-DIT-ARM-REACH` closes it with only the DiT half.** Deleting the two-line call leaves `test_minimax_music3_ar` 37/37 · 640/640 and `test_minimax_music3_speech` 9/9 · 223/223 green ([`minimax-music3.md`](../specs/minimax-music3.md) §19.5 carries the mutation and the binary hashes). Two things stop an existing gate from seeing it, and the SECOND is the one that matters: `--speech-device 1` is refused by name on a CPU-only build before a queue exists, AND §19.6's "device path TAKEN" leg rides `test_minimax_music3_ar`, whose observable `Music3DepthDeviceForwardCount()` is a counter §19.5 itself records as unreachable from production — its only readers are the tests written for it (`test_minimax_music3_ar.cpp:1325,1351,1583,1589,1753,1758`). The instrument that WOULD answer the call site is `ar.depth_staging`, emitted at `minimax_music3_llm.cpp:582` and read by nothing. NOT FIXED IN FLOW and the reason is precise: closing it needs the shipped engine on a real accelerator against the 28.5 GB checkpoint inside an `rc` lease, which is a second GPU leg and a second gate file, not a repair to the row in flight. It is closable by exactly `MUSIC3-DIT-ARM-REACH`'s method — a `gpu;checkpoint;music3`-labelled parity gate entering through `include/vllm.h` with `device = 1`, exiting 77 without its preconditions, asserting `ar.depth_staging` `calls == 1` with the host bucket absent — and that row's `thor:gpu0` job `f63f60e8-957a-4062-92f8-54e5bbb49d92` already FIRED `ar.depth_staging` once without asserting it, so the instrument is known live on the real path. Owed under `## Owed` in [`minimax-music3.md`](../specs/minimax-music3.md) §19.7 | bug | -| [#1843](https://github.com/mudler/vllm.cpp/issues/1843) | `FIX-FP8-PLAN-CAPTURE-1843` | The fp8 cuBLASLt lane still queries `cublasLtMatmulAlgoGetHeuristic` inside CUDA-graph capture because `VT_FP8_PLAN_CACHE` ships OFF (`fp8_plan_cache.h:49-59` @ `364f2a898`), so a captured decode on an fp8-tower model dies on CUDA 13.3 even with PR #1741 -- measured on `dgx:gpu0` (GB10, staged CUDA 13.3.73): #1741 alone fails on the fp8 lane, `VT_FP8_PLAN_CACHE=1` alone fails on the bf16-TN lane, both together pass the graphed 35B gate token-exact on all three arms. The fix is the default flip #1741's spec owed, with the same not-a-performance-knob polarity argument its `gemm_plan_cache.h` records. Claimed by row `FIX-FP8-PLAN-CAPTURE-1843` ([spec](../specs/fix-fp8-plan-capture.md)) | bug | -| [#1849](https://github.com/mudler/vllm.cpp/issues/1849) | `SPEC-DFLASH2` | **The DFlash2 draft step costs a flat ~23 ms at EVERY K, and the two levers #1849 names resolve differently once read from the records.** Lever A (quantize the shared head) is ALREADY LANDED for the measured subject: `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` stores `lm_head` as W4A16_NVFP4 g16 (header-verified 2026-08-21, quantization-matrix `QUANT-QWEN38-27B-NVFP4-ARM`), upstream computes with it packed through `lm_head.quant_method.apply`, and both our reads have been packed since #1628 — so the head traffic is ~2×0.72 GB not 2×2.54, the draft-phase floor re-derives to ~9 ms, and the unattributed residual GROWS to ~13-14 ms. Lever B (launch/sync trim) is counted in code at one replay + ~10 launches + ~76 B up / 64 B down + one sync — well under 0.5 ms, so the residual sits INSIDE kernels and needs on-box attribution. W9 lands `VT_SPEC_TRACE=2` (the `[spec-phase-dev]` pre/fwd/select/walk split) as the instrument, and borrow-first loading for the draft's shared bf16 embed+head (~5.1 GB host on the bf16 arm, ~2.5 GB on the r0b0tlab arm; memory only, no step-time claim). The bf16-target arm's 2×2.54 GB head reads are upstream's own serving dtype and stand as a recorded ceiling. Wave spec [dflash2-draft-fixed-cost.md](../specs/dflash2-draft-fixed-cost.md); the K-ladder rerun, the `ncu`/`nsys` attribution and any step delta are owed there, operator-run | perf | -| [#1844](https://github.com/mudler/vllm.cpp/issues/1844) | `ENG-MM-INPUT-PIPELINE` | **`scripts/mm/tower_skip_rss.sh` killed every measured leg mid-load, because `run_arm`'s `/health` poll was answered by the PREVIOUS leg's server.** First real run (`thor:gpu0`, worker `rc-worker-kk96r`, `d60692c8`): checkpoint staged and verified (29 files, 8887294190 B), both binaries built sha256-identical (`78d582e4...`), live target query green on both build dirs -- and then **five 0-byte `.time` files** and `VOID` on both pairs. `warmup` reached `listening on http://0.0.0.0:18607` (1286 B log); all four measured legs stopped at `loading model from ...` (363/345 B) inside one minute. TWO defects, one shape. (a) `$PORT` is fixed and the poll starts immediately, so a stale listener answers it and the leg is ready before it has read a tensor. (b) `kill "$pid"` signals `/usr/bin/time`, which installs no handler: the timer dies before writing its `-o` file and the server is reparented to init and KEEPS THE PORT -- which is what was answering. Measured: `/usr/bin/time -v -o f sleep 100 & kill $!` leaves `f` at 0 B and `sleep` alive with ppid 1; signalling the CHILD leaves `f` at 752 B with a `Maximum resident set size` line. **Nothing could catch it**: `test_tower_skip_rss_report.py` was 60/60 green over finished files, and `run_arm`, the poll and the teardown ran only under a lease -- the residual [#1819](https://github.com/mudler/vllm.cpp/issues/1819) recorded verbatim, now observed. FIXED IN FLOW: a leg refuses to start into an occupied port; readiness requires the banner in the leg's OWN log (its stdout, which no other server can write) before `/health`; SIGTERM goes to the SERVER so the timer survives to write; the port must stop accepting before the next leg; each of the three waits is bounded and each bound REFUSES; and a leg whose `.time` carries no `Maximum resident set size` line fails AT that leg rather than as VOID four legs later. GATED: `TOWER_SKIP_RSS_SOURCE_ONLY=1` sources the harness for its functions alone and `tests/scripts/test_tower_skip_rss_arm.py` drives `run_arm` against a fake server on a scratch port -- stale listener, absent banner, death during load, the happy path's non-empty `.time`, five legs in the declared order, and both halves restored as mutations. RED-first: 11 of 14 cases fail against the pre-fix `run_arm`, the stale-listener case reporting `LEG default OK` with a 0-byte `.time` | bug | -| [#2065](https://github.com/mudler/vllm.cpp/issues/2065) | `KV-FP8` | The ROCm fp8-e4m3 KV cache arm: the store kernel (`vt::ReshapeAndCacheFp8` for `DeviceType::kROCM`), the fp8 dequant on the ROCm paged-attention read, the `OpId::kReshapeAndCacheFp8` registration for `kROCM`, and the widening of the `src/vt/ops.cpp` fp8 read refusal to admit `kROCM`. Mirrors the CUDA W2 arm element-for-element; the CPU kernels remain the oracle. Spec [`fp8-kv-cache.md`](../specs/fp8-kv-cache.md) `## W6` | feature | -| [#1846](https://github.com/mudler/vllm.cpp/issues/1846) | `MODEL-MM-dots3-note-dots3-note-for-causal-lm` | **The released `dots3-note-prev` shard index declares `indexer_rope_layout: "leading"` and `indexer_rope_converted_from: "tail"` in its `metadata` block, and NOTHING reads either key** — `git grep indexer_rope_layout` over vLLM `origin/main` returns nothing. Measured at W2 while reading the whole index. It is the publisher stating how the DSA indexer's `wq_b`/`wk` are laid out along the 128-wide index head, and it agrees with what upstream's code does anyway: `DeepseekV2Indexer` rotates `[..., :rope_dim]` and leaves `[..., rope_dim:]` (`deepseek_v2.py:805,:814`, `rope_dim` 64 of `index_head_dim` 128), which is a LEADING slice. NOT spec §4 trap 2: that one is about which PAIRS rope rotates (GPT-J vs NeoX), this one about which HALF of the head it rotates, and both are numerically silent on a row spec §6.4 says has no oracle. W2 pins both values in an assertion so a re-published checkpoint cannot flip the layout silently; W2 consumes neither, because W2 writes no maths. W3 owes the slice | feature | -| [#1857](https://github.com/mudler/vllm.cpp/issues/1857) | `SPEC-DFLASH2` | **The q>1 DFlash2 verify rides the PREFILL flash lane (`is_prefill = num_tokens > num_reqs`, `fa2_decode` hard-requires `num_tokens == num_reqs`), costing +9 ms/step from q=2 to q=9 on the #1574 K-ladder — the last attributed gap against SGLang (27.60 vs 25.07 tok/s at equal acceptance, ~109 vs ~122 ms/step, both paying the same ~20 ms draft).** W10 mirrors upstream's spec-as-decode: the reorder-threshold policy `1 + (parallel_drafting ? 2 : 1) * K` (`backend.py:718-736` @ `b389ac2946`, identical at the pin) classifies the runner's already-verified uniform verify length onto the decode class, the classification travels `CommonAttentionMetadata -> PagedAttentionArgs`, and a new ADDITIVE d256 launcher serves it with the exact presentation upstream `mha_fwd_kvcache` uses at seqlen_q>1 — batched split-KV, bottom-right causal against `seqused_k` (the draft mask with no new mask code), `set_params_splitkv` heuristic. The shipped q==1 arms and every unclassified batch stay dispatch-identical; `VT_FA2_SPEC_DECODE=0` restores the prefill route for a same-binary A/B. Wave spec [dflash2-spec-as-decode.md](../specs/dflash2-spec-as-decode.md); the GPU step-time delta (the −8-9 ms claim), the GPU token gates and the first CUDA compile are owed there, operator-run | perf | -| [#1853](https://github.com/mudler/vllm.cpp/issues/1853) | `LTX25-DIT-ATTN-FLASH` | **`PENDING` on a `dgx:gpu0` lease: the arithmetic-perturbation reference render that would make [#1743](https://github.com/mudler/vllm.cpp/issues/1743)'s criterion RELATIVE.** §11 of [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) relocates the pixel verdict onto **correspondence** and **incoherence**, which discriminate a degraded render from a separated trajectory. It deliberately does NOT answer the other half of #1743: is the swap's divergence no worse than this pipeline's own divergence under an arithmetic perturbation of comparable size. That needs one further arm - the **naive** path at `768x448/49f`, seed `20260820`, on §10.7's pinned binary and checkpoints, with a bounded `+/-1` bf16 ULP dither injected at the DiT attention output at the `8.6e-05` to `3.7e-04` per-element flip rate §10.2 derives - after which `D(flash, naive) <= D(dither, naive)` is a bound with NO chosen constant. **No lease was authorised for #1743, so this is PENDING and not skipped.** The cross-build `baseline-20260820` vs `naive` figure (mean \|delta\| **9.452407**, LARGER than the swap's **6.414156**) is NOT that control and is not used as one: the binary lineage differs, so every other commit between `a50c57d69` and `3e2961ef0` sits inside it, which §10.8 already records. NOT FIXED IN FLOW: it needs a GPU lease this work does not have. Owned by row `LTX25-DIT-ATTN-FLASH` and listed under `## Owed` | bug | -| [#1854](https://github.com/mudler/vllm.cpp/issues/1854) | `LTX25-DIT-ATTN-FLASH` | **NOT GATEABLE in this tree, declared rather than proxied: nothing asks whether an LTX-2.5 render is GOOD, only whether two renders are the SAME.** Raised while relocating #1743's criterion. The RELATIVE form of the question IS now answered and gated - the coherence checks of §11.3 assert that neither arm is systematically sharper, blockier, quieter or less mobile than the other, at `K <= 0.5` where a one-directional degradation gives `K = 1` exactly. The ABSOLUTE form is not answered: **prompt adherence needs a vision-language model**, which §10.8 already refuses to approximate ("a check for 'is this a golden retriever shaking off water' is a model, not a threshold"), and **artefact-freedom needs an absolute reference render** from an oracle that runs this pipeline, which `.agents/oracles/` does not have. `scripts/ltx25-render-compare.py` therefore computes an **absolute quality panel per arm** - 8-grid and 32-grid blockiness ratios, clipped-pixel fraction, mean sharpness - prints it, records it in the JSON and **checks none of it**, saying so in its own output, rather than inventing a threshold that means nothing without a reference. NOT FIXED IN FLOW and deliberately not: a proxy for perceptual quality is the `a-shape-valid-gate-passes-a-wrong-artefact` failure. Owned by row `LTX25-DIT-ATTN-FLASH` and listed under `## Owed` in §11.5 | bug | -| [#1855](https://github.com/mudler/vllm.cpp/issues/1855) | `LTX25-DIT-ATTN-FLASH` | **The DiT attention swap leaves the VIDEO directionless and takes 4% off the AUDIO in the only passage that has any.** Found by the structural criterion of [#1743](https://github.com/mudler/vllm.cpp/issues/1743) on the frames [#1612](https://github.com/mudler/vllm.cpp/issues/1612) had already rendered, with NO GPU and no new lease. `flash` vs `naive`, one binary `834cec55...`, `768x448/49f`, seed `20260820`. **Correspondence all passes**: worst frame margin **1.4230** (must exceed 1), **0 of 49** frames matching better at any spatial offset other than `(0, 0)`, audio argmax at lag **0**. **The three VIDEO coherence statistics are incoherent**, an order of magnitude above their own `N^-1/2` floor and two to three orders below a direction, with the majority sign a coin: sharpness `K` **0.032512** (N 65856, floor 0.0039, sign 0.502), blockiness **0.007914** (N 7448, floor 0.0116, sign 0.501), motion **0.031635** (N 64512, floor 0.0039, sign 0.495). **The AUDIO is not**: `K` **0.674002** over 376 windows against a floor of 0.0516, RMS ratio `flash/naive` **0.962289** - the 3.3% §10.7 printed and never checked. It is CONCENTRATED, not a gain change: the 2.005 s track is near-silent outside one passage and the whole effect is a **4.0% loss in windows 125-249** (2471.7 vs 2573.2) with the near-silent thirds at 1.0116 and 0.9935; 182 windows quieter and 194 louder, losses **5.1x** the gains in magnitude. **TWO LIMITS, stated**: 376 windows is NOT 376 independent observations because the track has ONE loud event, so the audio verdict rests on a single acoustic passage while each video verdict rests on tens of thousands of tiles; and `K = 0.674` is a PARTIAL direction sitting between the floor and 1, so the `0.5` constant IS load-bearing there and any constant above 0.674 would not fire. **The cross-build pair is the CONTROL-shaped context and it goes the other way**: `baseline-20260820` (`a50c57d69`, ancestor, naive path) vs today's naive reads audio `K` **0.312163** and RMS ratio 0.986061, and reads `SEPARATED, NOT DEGRADED` overall despite a LARGER mean \|delta\| (9.452407 vs 6.414156) - which is the ordering a widened tolerance cannot produce. **NOT the shipped default**: these are the `flash` rung of #1549, and [#1551](https://github.com/mudler/vllm.cpp/issues/1551)'s FA-2 arm, whose unset default is what `main` builds today, has never been rendered at production geometry. **NOT attributed**: why a reassociated attention sum costs 4% of audio amplitude while leaving the picture directionless is unexplained. NOT FIXED IN FLOW: a finding about a change already on `main`. Recorded in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) §11.8 | bug | -| [#1863](https://github.com/mudler/vllm.cpp/issues/1863) | `BACKEND-ROCM` | **ROCm decode loses 2.4x to llama.cpp on a DENSE model, so the deficit is not MoE-specific.** `Ornith-1.5-9B-Q4_K_M.gguf` (arch `qwen35`, 5.23 GiB, 8.95 B, no MoE block, no router, no grouped expert GEMM, no offload) on RX 9060 XT (gfx1200), ROCm 7.2.3, batch 1, 32 tokens: vllm.cpp `4b1154bc5` **18.393 / 18.574 t/s** against llama.cpp `b10451` (`10bf611e5`) HIP **44.24 +/- 6.10 t/s**. The 14B MoE measured in the same session for context is 13.000 / 13.145 against 38.06 +/- 8.01, a 2.90x gap, so the dense and MoE deficits are the same order. This REFUTES scoping the ROCm decode gap to expert routing or to `kMatmulBTQuantGrouped`, which is where [#1400](https://github.com/mudler/vllm.cpp/issues/1400) and [#1294](https://github.com/mudler/vllm.cpp/issues/1294) left it: a model that touches none of that still loses 2.4x. Named hypothesis, NOT tested: `QuantizeQ8KK` has a second call site at `rocm_grouped_gemm.hip:479` inside `MatmulBTQuantKernelRocm`, the non-grouped dense path that a dense k-quant GGUF now reaches because [#523](https://github.com/mudler/vllm.cpp/pull/523) registered `kMatmulBTQuant` on ROCm; at dense decode its launch is `m * nsb` with `m` = 1, thinner than the `m` = T x top_k case #1400 profiled at 35% of GPU time. One parallelization fix would move both paths IF the quantizer is the cause, and nothing here establishes that — a `rocprofv3 --kernel-trace` pass on the 9B is what would. Stated rather than implied: the two sides are not measured with one construction (`vllm-cli` divides completion tokens by whole-call wall time, `llama-bench` reports `tg32`; #1400's differencing method is what a gate on this number owes), error bars are 14% and 21% at 3 repetitions, the host was not idle at loadavg 2.2-3.0 although free VRAM was asserted above 13 GiB with no resident model process, and this is one prompt, one quantization, batch 1 | perf | -| [#1865](https://github.com/mudler/vllm.cpp/issues/1865) | `SPEC-DFLASH2` | **W10's spec-as-decode lane never engages at runtime: the q=9 verify still runs `PagedFlashKernel` (nsys, 16 x 117 calls — the real prefill forward included) and `LaunchSpecDecodeFA2Bf16` never appears, with the on/off A/B speed-neutral.** Traced: the threading (runner -> `CommonAttentionMetadata` -> `PagedAttentionArgs`) is live — a probe at `vt::PagedAttention` under the production CPU fixture shows every uniform verify arriving classified — but the W10 admission dies on its bf16-query conjuncts because the model-side dtype selection (`FullAttnBlockPaged`) has NO spec-as-decode arm: the verify's bf16-ness rides the PREFILL lever's `Fa2PrefillOn()`, while the CUDA admission reads `Fa2SpecDecodeEnabled()`/`Fa2Decode*Enabled()` — two sides consulting different switches, and the profiled binary's FA2 arm was dark end to end (unstaged-CUTLASS configure prints `CUDA FA2 compiled-arch manifest: []` and builds green). Repair: the eligibility is extracted to a host-testable seam (`ClassifyDenseFa2`) and gains the spec arm reading the SPEC lane's own toggles; `vt::PagedAttention` counts classified arrivals so the W10 review's dead mutation (deleting the `pa_args.uniform_spec_query_len` threading) reds on a CPU box; a classified batch the CUDA dispatch cannot serve narrates ONCE to stderr naming the failed conjunct group. The nsys re-profile on a manifest-verified FA2 build and the moving A/B are owed, operator-run — wave spec [dflash2-spec-as-decode-repair.md](../specs/dflash2-spec-as-decode-repair.md) | bug | -| [#1877](https://github.com/mudler/vllm.cpp/issues/1877) | — | NVFP4 35B greedy decode is not run-to-run reproducible at c1 for requests after the first: two identical invocations (same binary, env, seed 777, temp 0) diverge at tokens 21 and 457 on requests 2-3, while request 1 is byte-identical and the bf16 35B null is clean on the same harness. Sibling of #1283 (which was c16); candidates: marlin/grouped-MoE atomics, per-process fp8 plan selection, cross-request state. Makes arm-identity-at-depth undecidable on this checkpoint. Found by the GDN-MOE-PACKED-BA speed A/B; listed under `## Owed` in [gdn-moe-packed-ba.md](../specs/gdn-moe-packed-ba.md) | bug | -| [#1878](https://github.com/mudler/vllm.cpp/issues/1878) | — | Packed GDN decode diverges from the rollback arm within 950 greedy tokens on the bf16 35B (first difference token 33 request 2, token 213 request 3, reproduced exactly across pairs) with a CLEAN same-arm null -- a deterministic kernel-numerics difference (FLA cubin bf16 vs split F32 pair), the 27B near-tie class, recorded as a measured property of the lever pending a quality-at-depth or oracle-continuation disposition. Found by the GDN-MOE-PACKED-BA speed A/B; listed under `## Owed` in [gdn-moe-packed-ba.md](../specs/gdn-moe-packed-ba.md) | bug | -| [#1890](https://github.com/mudler/vllm.cpp/issues/1890) | `SPEC-DFLASH2` | **The DFlash draft block's attention never got the FA-2 treatment: `DFlashPagedBlockAttentionWarpKernel` at 449.7 us/call x 5.1 calls/step = 2.29 ms/step, against SGLang's 14.3 us/call for the same work** — while W10's TARGET VERIFY, on the FA-2 split-KV lane, runs at 17.1 us/call and is marginally ours. 16.2 verify + 5.1 draft = 21.3 attention calls/step on both engines: same call count, one lane 31x slower. **The blocking property is KV RESIDENCY, and it is none of the four #1890 names** (page size 16, identity block table, bf16, head dim — all admit): `vt::DFlashPagedBlockAttention` reads the block's own (1+k) K/V out of contiguous per-layer tensors that are in NO paged cache, and every split-KV launcher addresses K and V exclusively through a block table, so handing one the store's pools would drop every block row — a wrong answer, not a slow one. W11 makes those rows RESIDENT (`vt::ReshapeAndCache` into the store's own pages) and reads the whole combined sequence as ONE `vt::PagedAttention`, which is the presentation upstream uses for the same work and is why SGLang issues one kernel for both lanes. The mask maps exactly with no new mask code — `PagedAttentionArgs` already carries FlashAttention's bottom-right alignment, so full / SWA / plain-causal become non-causal / causal+`{W-1,0}` / causal. **BYTE-IDENTICAL on CPU, asserted rather than argued**: the two kernels are the same three-pass online softmax in the same j-ascending order over the same bf16 bits, gated element-for-element across five mask and layout cases plus a drafted-token A/B through the production runner. CUDA additionally widens the W10 admission and `LaunchSpecDecodeFA2Bf16` to head dim 128 and to the three masks (the d256 verify arm stays dispatch-identical), and a classified batch the dispatch cannot serve now NARRATES its failed conjunct group once to stderr. `VT_FA2_DFLASH_BLOCK=0` is the same-binary rollback. The GPU number, the on/off A/B, the GPU token battery and the first CUDA compile are owed, operator-run, in [dflash2-draft-block-fa2.md](../specs/dflash2-draft-block-fa2.md) | perf | -| [#1894](https://github.com/mudler/vllm.cpp/issues/1894) | `SPEC-DFLASH2` | **The DFlash2 runner fixture drafts a CONSTANT — `19 19 19` at all eight steps — so every drafted-token comparison through it is a tautology against a numerics change.** Found by #1890's mutation pass rather than by reading, with a one-off probe that printed what `DraftedBlocks` returns. Three mutations of the draft-block attention, each a genuinely WRONG attention (mask polarity forced causal; the paged K/V write neutralised; the read handed the store's `seq_lens` instead of the extended bound), left the comparison GREEN while the byte-for-byte op gate `test_qwen3_dflash_block_route` red on all of them — which is how the degeneracy was located. **Not only a W11 concern**: the landed `dflash2 runner (W8): the paged lane and the materialized lane draft identically` case compares the same constants, so its stated guarantee is not measured by it either. NOT degenerate under every perturbation — the D9-scalars case does observe a difference — so the fix is weights whose per-position argmax actually separates, applied without disturbing the W3/W4/W9/W10 cases that read the same fixture. NOT FIXED IN FLOW: changing the fixture's weights moves five landed cases at once and needs its own red-before evidence, a different unit of work from #1890. Owned by row `SPEC-DFLASH2` and listed under `## Owed` in [dflash2-draft-block-fa2.md](../specs/dflash2-draft-block-fa2.md) | bug | -| [#1866](https://github.com/mudler/vllm.cpp/issues/1866) | `KERNEL-GEMM-FP8` | **The FP8 tower's +3.04 ms/step is not an autotuner's absence: vLLM's fp8 GEMM is a CUTLASS M ladder we ported and truncated.** #1866 read #1857's corrected head-to-head as "our cuBLASLt tower runs `sm89_xmma` 32x64x64 where SGLang runs `nvjet_sm121`, and their server sweeps `fp8_gemm` at startup while we do none". Read at the pin `5559679229`, vLLM does neither: its CUDA fp8 backend order is Marlin -> FlashInfer -> Cutlass -> PerTensorTorch (`vllm/model_executor/kernels/linear/__init__.py:325-334`), a Cutlass-capable device takes `ops.cutlass_scaled_mm` (`.../scaled_mm/cutlass.py:265`), `git grep "cublasLt\|AlgoGetHeuristic" -- csrc vllm` is EMPTY, and the only startup sweep tunes FlashInfer ops plus a bf16 router GEMM (`vllm/model_executor/warmup/kernel_warmup.py:47-72,189,238-242`). Its fp8 small-M behaviour is a STATIC four-way tile ladder (`csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh:155-176`, reached on GB10 via `scaled_mm_entry.cu:222-225`'s `version_num >= 120`): `M<=16` 16x64x128 EpilogueTile 16x32, `M<=32` 32x64x128 EpilogueTile 32x32, `M<=256` 64x64x128, else 128x128x128. **We ported that file and kept two rungs**, recording the other two as "perf-only for tiny M and covered correctly (predicated) by the M64 pingpong tile" — true, and tiny M IS decode, so a batch-1 step computed a 64-row tile for one row and the 9-row spec-decode verify computed one for nine. A wrong tile is a SLOW answer, which is why every token gate stayed green over it. The two rungs are restored behind `VT_FP8_CUTLASS_SMALL_M` with the ladder itself in a CUDA-free header so a host with no `nvcc` gates it by value; the record that called our cuBLASLt lane "the native equivalent of vLLM's nvjet_sm121_qqtst kernels" is corrected in the same change. **STAYS OPEN**: the default arm is still cuBLASLt (`VT_DENSE_CUBLASLT_FP8`), so a production decode step does not reach the new rungs, and the CUDA compile, the token gates and the `VT_DENSE_CUBLASLT_FP8` decode A/B this finally makes fair are all listed under `## Owed` in [perf-fp8-small-m-dispatch.md](../specs/perf-fp8-small-m-dispatch.md). No speed number is claimed | perf | -| [#1883](https://github.com/mudler/vllm.cpp/issues/1883) | `TEST-REG-SERVER-GUARD` | `test_minimax_music3_e2e_real` was registered outside `if(VLLM_CPP_SERVER)` while calling `ApiServer`, whose translation unit `CMakeLists.txt` compiles only inside it, so every `-DVLLM_CPP_SERVER=OFF` configure succeeded and then failed at `ld` with 16 `undefined reference` lines -- the standing red on `main`'s `build-test-cpu-arm64-full` job. Fixed in flow: the registration moves inside the guard beside the other server-linking suites, and `scripts/check-test-registration.py` gains a static `server_guard_errors` check that derives the gated translation units from the top-level CMake, resolves their declaring headers, walks the tree's own include graph transitively, and refuses any `vllm_cpp_add_test` target outside the guard that reaches one. Spec [test-registration-server-guard.md](../specs/test-registration-server-guard.md) | bug | -| [#1892](https://github.com/mudler/vllm.cpp/issues/1892) | `TEST-REG-SERVER-GUARD` | `check-pr-size.py`'s checker-change evidence harness ran `tests.scripts.test_check_test_registration` with neither `cmake`, `ctest` nor `ninja` reachable: `EVIDENCE_REQUIRED_TOOLS` named only the windows-portability module, so the sanitized `PATH` (`os.defpath` plus an empty private tools directory) could not start the programs that module drives. CI reported `FileNotFoundError: 'cmake'` with 22-26 errors, all charged to the checker under change rather than to the harness -- the broken-instrument shape, and the same gap [#458](https://github.com/mudler/vllm.cpp/issues/458) closed for the other module. Invisible until now because the harness only runs when a checker and its evidence file change together. Fixed in flow with #1883: the module declares its tools, and the test DERIVES the expectation from the checker's own argument-list literals rather than transcribing a list, so the `ctest` that the first fix missed reds locally instead of in CI. Spec [test-registration-server-guard.md](../specs/test-registration-server-guard.md) §10 | bug | -| [#1875](https://github.com/mudler/vllm.cpp/issues/1875) | `MODEL-DSV4-EXL3` | MiaAI-Lab/DeepSeek-v4-Flash-One-DGX-Spark runs a REAP-pruned K216 DeepSeek-V4-Flash at a claimed 44-47 tok/s decode on one GB10 -- EXL3 3.0bpw trellis quant (~99.5 GiB, the first V4-Flash quant that fits one Spark), SparkInfer (NVIDIA-vLLM-26.02 fork) with K5 speculative decoding. Developer direction: load the same quants and match or beat the speed. vLLM has no EXL3 at the pin, so the row proposes `exllamav3` @ `2398c056` as a pinned secondary oracle (its HEAD carries DSV4 support). Spike 2026-08-24 on the issue pins the full format (MCG `0xCBAC1FED` 3-instruction decode, 16x16 tail-biting trellis tiles, H128+sign vectors, no scales, lossless TP4-to-TP1 coalescing, K216 physical compaction our config-driven loader accepts unchanged). Caveats recorded: their number includes spec decode with no bare-AR figure, and the checkpoint's own README says end-to-end generation is runtime_pending. Row [spec](../specs/model-dsv4-exl3.md) | feature | -| [#1864](https://github.com/mudler/vllm.cpp/issues/1864) | `ENG-UPSTREAM-LTX2-PIN` | `Lightricks/LTX-2` is now pinned at `fd4ded7f2d88d3da713abcdd4ad41ecc4a9314ca` in [`.agents/oracles/ltx-2.md`](../oracles/ltx-2.md), and it records `gateable = no` because the oracle has never RUN the model — this issue is what that field owes. AGENTS.md sets the bar at "demonstrably builds and runs the model. Constructing a config proves nothing." MEASURED: thirteen tracked scripts import and execute upstream `ltx_core` / `ltx_pipelines` code, and every one runs individual MODULES at reduced dimensions on synthetic PRNG weights, or reads constants and safetensors headers. The committed goldens say so in their own headers — `tests/vllm/models/ltx2_vae_goldens.inc:3-6`: "Weights and inputs come from the shared deterministic stream, so no weight byte is checked in" — and the same sentence appears in `ltx2_goldens.inc`, `ltx2_pipeline_goldens.inc`, `ltx2_text_goldens.inc`, `ltx2_tiling_goldens.inc` and `tests/vllm/multimodal/ltx2_image_cond_goldens.inc`. Exactly two scripts touch real checkpoint bytes and neither runs the model: `scripts/measure-ltx2-prompt-adaln.py:124-133` forwards ONE `AdaLayerNormSingle` plus the per-block tables out of a 21 B DiT, and `scripts/measure-ltx2-keyframes-meta.py:156,203` builds the model on the META device and runs upstream's loader with no forward pass in the file; neither writes a committed artifact. Negatives with their searches: `ls tools/oracle/` holds `music3_oracle.py` and `README.md`, `grep -rn -i ltx tools/` returns no hit across the 72 files `git ls-files tools` reports, and none of the 103 entries under `tests/parity/goldens/` matches `ltx` case-insensitively. A second finding the run would close: where the revision IS asserted (`scripts/gen-ltx2-res2s-goldens.py:397-402`, a hard `SystemExit`) no weight is loaded, and where weights are loaded (`scripts/measure-ltx2-prompt-adaln.py:60-66`) only the interpreter PATH is asserted, never the revision — and `.agents/specs/ltx-2-5.md` §7.0(b) records a decoy `ltx_core` that produced byte-identical goldens and exited 0. NOT fixed in flow and deliberately: it needs a GPU lease and the gated `Lightricks/LTX-2.5` checkpoint, and no lease was authorised for the filing row. Owned by `ENG-UPSTREAM-LTX2-PIN` and listed under `## Owed` in [oracle-ltx-2-pin.md](../specs/oracle-ltx-2-pin.md). Does NOT duplicate [#633](https://github.com/mudler/vllm.cpp/issues/633) (the vLLM-Omni pin, a different repository) or [#1854](https://github.com/mudler/vllm.cpp/issues/1854) (an absolute render-quality gate, which states the same absence from the other end) | bug | -| [#1885](https://github.com/mudler/vllm.cpp/issues/1885) | `LTX25-DISTILLED-LORA-REQUIRED` | **`test_ltx2_video` reports a DIFFERENT assertion count on every run of the same binary, so any count comparison across a diff measures noise.** Seven runs across two binaries at `5c789015e` (Release, x86-64, `SERVER=ON`) gave 4337, 4338, 4339, 4340, 4341, 4342 and 4343 — a 7-assertion spread — while `test cases: 105 \| 105 passed \| 0 failed \| 0 skipped` held on every one. The two binaries are two worktrees at the same commit, so it is neither a source nor a configuration difference; [#1445](https://github.com/mudler/vllm.cpp/issues/1445)'s spec first explained it as `SERVER=OFF` against `SERVER=ON` and that explanation is wrong. It MATTERS rather than being cosmetic because this repository quotes assertion counts as evidence — `docs/FEATURES.md` carries `test_ltx2_dfr` 11/11, 652 assertions and `test_ltx2_tiling` 10/10, 915 assertions — and a reviewer comparing 4343 before against 4337 after will hunt for six assertions a diff deleted when it deleted none, which is a failure that reads as a result. It also puts a genuine five-assertion deletion inside the observed noise. NOT ESTABLISHED and deliberately not guessed at in the record: WHICH cases vary (timing-dependent loops in the phase-log live-tick and RSS-style cases are the obvious candidates and nothing has isolated them), and whether any other suite shares the property — `test_ltx2_pipeline` returned 3475 on three runs, which is three runs and not a claim. Owned by row `LTX25-DISTILLED-LORA-REQUIRED` and listed under `## Owed` in [ltx25-distilled-lora-required.md](../specs/ltx25-distilled-lora-required.md). Not fixed in flow: the varying cases are pre-existing, that row touches none of them, and isolating them needs a bisection over 105 cases | bug | -| [#1872](https://github.com/mudler/vllm.cpp/issues/1872) | `LTX25-DIT-ATTN-FLASH` | **`align.audio_lag` is a bare argmax with NO margin, and on real frames it fires on a correlation difference of `4.5e-05`.** Found while measuring [#1855](https://github.com/mudler/vllm.cpp/issues/1855)'s audio direction on the frames [#1612](https://github.com/mudler/vllm.cpp/issues/1612) rendered, with no GPU and no lease. `flash` against `baseline-20260820` reads `best lag -1 samples` on `r = 0.926353` there against `0.926308` at lag 0 - **one sample of 96,480 at 48 kHz, 20.8 microseconds** - and the whole run reads `READING MISALIGNED` and exits 1 on it, which §11.6 defines as the state where the pair is not comparable at all. **The asymmetry is the defect.** `align.frames` was written WITH a bound - `margin > 1`, computed by `frame_correspondence` and printed in the report - precisely because "nearest" without "by how much" is not a correspondence. `audio_correspondence` computes the same shape of quantity, returns `best_lag` plus `r` at the best lag AND at 0, and the check then compares only the integer, over a `+/-2000` sweep of 4001 float candidates with no tie handling. **NOT a case for widening a threshold** (#1668 and §9 forbid that): the repair ADDS a bound that does not exist rather than moving one that does. **NOT REPAIRED IN FLOW**, because it changes checker semantics and therefore owes its own row, spec section, a red-before test on both a genuinely shifted track and a hairline tie, and a fresh review. **No published verdict moves**: the same-binary `flash` vs `naive` pair of §10.7 reads `best lag 0` and passes, so #1743 and #1855 stand. Owned by row `LTX25-DIT-ATTN-FLASH` and listed under `## Owed` | bug | -| [#1881](https://github.com/mudler/vllm.cpp/issues/1881) | `LTX25-DIT-ATTN-FLASH` | **The LTX-2.5 pixel A/B recorded a 75 KB LAUNCHER as its binary identity, and two builds a whole release window apart printed the SAME value while the 92 MB library that holds every measured kernel differed by 6,372,624 bytes.** Observed live on `dgx:gpu0` 2026-08-24 while the FA-2 ladder of [#1855](https://github.com/mudler/vllm.cpp/issues/1855) was building. `BINSHA=$(sha256sum "$BIN/ltx2-gen")` hashes 75,344 bytes of `main()`; `vt::Attention`, `vt::AttentionDenseFlash`, `vt::AttentionDenseFa2`, both VAEs and the loader are all in `libvllm.so.0.0.3`, which was hashed NOWHERE. Run `1612-r3` (source `3e2961ef0`, `libvllm` 85,703,328 B) and run `1853-fa2-r1` (source `62cbae10d`, `libvllm` `f046e75dcede2586...`, 92,075,952 B) both recorded `binary_sha256=834cec557c16cf77...`, each `binary_built=in-lease` with `BUILD_RC=0`. The launcher's own translation unit did not change, so its output is reproducible BY CONSTRUCTION - the one artefact whose hash was stable was the one containing none of the code under measurement. Consequence: §10.7's "the binary `834cec55...`", which #1743 and #1855 rest on, does NOT pin the code that produced them, and a later reader reads the same string and concludes the same code ran. **NOT an invalidation of 1612-r3**: its four arms ran from one build in one lease and each proved its own op from its own log; what its RECORD cannot do is tell its build from a later one. **FIXED IN FLOW for the pixel harness** - `LIBSHA` is computed, printed, written to `PROVENANCE` as `library_sha256` and added to every arm's `render.log` header, with three tripwires in `test_ltx25_pixel_ab_harness.py` that each red when their site is deleted. **STILL OWED**: `ltx25-dit-attn-flash-ab.sh` and `ltx25-dit-attn-fa2-hd128-ab.sh` carry the identical idiom, so every speed number they have recorded has the same hole. Owned by row `LTX25-DIT-ATTN-FLASH` and listed under `## Owed` | bug | -| [#1886](https://github.com/mudler/vllm.cpp/issues/1886) | `LTX25-DIT-ATTN-FLASH` | **The FA-2 arm that every production render resolves reads `DIRECTIONAL` against `naive` on the audio, and the checked statistic ranks it BACKWARDS against `flash`.** Measured on `dgx:gpu0`, `rc` job `4dcdd916...`, `RUN_ID=1853-fa2-r1`, source `62cbae10d`, `libvllm.so.0.0.3` `f046e75dcede2586...` (92,075,952 B), `768x448/49f`, seed `20260820`, FOUR arms from ONE build in ONE lease. This is the render [#1855](https://github.com/mudler/vllm.cpp/issues/1855) said had never been taken: `ltx2_device.cpp:536` is the only reader of `VLLM_LTX2_DIT_FLASH_ATTN`, unset selects `vt::AttentionDenseFa2`, and nothing in `include/`, `examples/` or `tools/` sets it. **Every arm proved its own op and neither other, and the counts were re-derived from each arm's raw `render.log` rather than read off its `ARM` summary**: `fa2` op18=0 op21=0 op22=1 (2.223 s/forward), `naive` op18=1 op21=0 op22=0 (45.512 s), `fa2-ctl` op18=0 op21=0 op22=1 (2.256 s), `flash` op18=0 op21=1 op22=0 (6.360 s). **The control is a zero, exactly**: `fa2-ctl` is BIT-IDENTICAL to `fa2`, 49/49 frames and `audio.wav` byte-equal, `control/treatment = 0.000000` on luma against a treatment effect of 8.952578, so nothing here is run-to-run variation. **Correspondence passes in full** (frame margin **1.1928 > 1** at frame 29, 0 of 49 frames off `(0,0)`, audio lag **0**) and **the three video statistics are incoherent** (sharpness **0.020738**, blockiness **0.040532**, motion **0.061146**). **The audio fires**: `coherence.audio_rms` `K` **0.511574** over 376 windows, means 865.774 vs 892.84, the shipped arm **3.03% quieter**, `top10% = +0.952`. `READING DIRECTIONAL`, `VERDICT FAIL (exit 1)`. **The direction is ONE CHANNEL**: ch0 `K` **0.705886** (-6.40%), ch1 `K` **0.060690** (-0.29%) at **1.18x** its own 0.0516 floor, i.e. FA-2 does not touch channel 1. **This CORRECTS #1855's attribution rather than confirming it**: `flash` loses the SAME channel (ch0 0.756589), so the direction belongs to the reassociated f32 online-softmax order both kernels share -- by TWO DIFFERENT reassociations, see `vt::AttentionDenseFast`, `vt::AttentionDenseFlash` and `vt::AttentionDenseFa2` in `include/vt/ops.h` and not to the #1549 swap. §11.9 pre-registered exactly this test before the render existed. **THE DILUTION IS ARM-DEPENDENT AND FALLS HARDEST ON THE SHIPPED ARM, which needed all three pairs to see**: the checked mono term retains 89.1% of `flash`'s channel-0 direction (0.674002 of 0.756589) and only 72.5% of `fa2`'s (0.511574 of 0.705886), because `flash` moves both channels the same way while `fa2` leaves ch1 at -0.29% and the average dilutes its channel-0 loss against an effectively unmoved channel. That is what pushed `fa2` to within 2.3% of the constant while `flash` sits 35% above it. **A first draft of this row claimed an INVERSION -- that the checked statistic ranks the two arms the wrong way round -- and a fresh review falsified it before it was published**: on `K` against `K` there is no inversion, because ch0 `K` ranks `fa2` below `flash` (0.705886 vs 0.756589) and the mono term ranks them the same way. The apparent reversal only appeared when an AMPLITUDE measure (`fa2` loses 6.40% of ch0 against `flash`'s 5.40%) was set against a COHERENCE measure, which is the category error §11.3 exists to prevent. The claim is withdrawn. **`fa2` vs `flash` shows no direction anywhere** (audio `K` 0.177718, sharpness 0.001954, `SEPARATED, NOT DEGRADED`, exit 0), which locates the direction between both fast kernels and `naive`. **The `flash` vs `naive` row is NOT an independent second observation**: it reads every §11.8 figure to six decimals because its renders are BYTE-IDENTICAL to run `1612-r3` across `3e2961ef0` -> `62cbae10d` and a `libvllm` differing by 6,372,624 bytes, so #1855 is reproduced rather than replicated and the count of independent observations is still ONE. **TWO LIMITS**: the checked constant carries the mono verdict by only 2.3% (`0.511574` vs `0.5`), a PARTIAL direction that §11.3's two-population argument does not cover, so there `0.5` is a chosen bound and NOT repaired by moving it (§9); and 376 windows are not 376 independent observations because the 2.010 s track has one loud passage. **NOT ATTRIBUTED**: why a reassociated attention sum costs 6.4% of one audio channel while leaving the other at its floor and the picture directionless is unexplained. NOT FIXED IN FLOW: a finding about a change already on `main` whose mechanism needs its own investigation. Owned by row `LTX25-DIT-ATTN-FLASH` and listed under `## Owed`. Recorded in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) §12.6 | bug | -| [#1887](https://github.com/mudler/vllm.cpp/issues/1887) | `LTX25-DIT-ATTN-FLASH` | **Three citations of `include/vt/ops.h:3304-3306` in [`ltx25-dit-attn-flash.md`](../specs/ltx25-dit-attn-flash.md) were stale and asserted the OPPOSITE of the sentence they supported.** Found by the fresh review of [#1871](https://github.com/mudler/vllm.cpp/pull/1871) while that change was adding a FOURTH citation of the same anchor. At `def85d285` those lines sit inside the `AttentionRelPos` doc comment and read "Reductions are strictly sequential per output element => thread-count independent and byte-reproducible", while the three sites (§4 numerics, §10.2, and §11.3 -- the ratified criterion section the whole pixel lane rests on) cite them for the claim that `vt::Attention` and `vt::AttentionDenseFlash` "differ only in association". Correct at `90e8c3c85`/`ff8f72807`, stale by `c4ba829a3`. **Correct anchors, each verified unique by phrase**: `ops.h:3315-3316` (`Fast` NOT bit-identical to `Attention`, different head_dim partial-sum grouping), `ops.h:3328-3329` (`Flash` order UNCHANGED from `Fast`, bit-identical), `ops.h:3381-3382` (`Fa2` NOT bit-identical to `Fast`/`Flash`, `mma.sync` reassociates QK^T and PV). **A second defect the same reading exposed**: the anchor was being used to claim `flash` and `fa2` share ONE reassociated order, and they do not -- they are TWO DIFFERENT reassociations, so the supportable claim is a shared CLASS, which is what §12.6 measures. FIXED IN FLOW in #1871: all three pre-existing citations repaired and the fourth corrected before publication. Owned by row `LTX25-DIT-ATTN-FLASH` | bug | -| [#1884](https://github.com/mudler/vllm.cpp/issues/1884) | `LTX25-PHASE-RESIDUE` | **`CheckSeamAnchor` proves a seam anchor's POSITION and not its MAGNITUDE, so a zero-width anchor passes it green.** Filed by the change that introduced the gate ([#1668](https://github.com/mudler/vllm.cpp/issues/1668)), against its own work, and MEASURED: mutation M7 -- `artifacts.mux` closed immediately after it opens, dropping the late `Close()`, so the render's whole tail is un-named again -- passes **796 of 796 assertions, Status: SUCCESS**. Every clause survives a zero-width window: the count is 1, it is neither `span` nor `nested`, it still opens after `artifacts.audio` ends, nothing follows it so it is still the last leaf, and the containment clause holds VACUOUSLY because a zero-width window has no interior, so every other leaf ends at or before it or starts at or after it. The same argument applies to `load.dit_config`. **Third appearance of one shape**, after the `decode.audio.mel` partial transfer the file already records ("an anchor covering half its call passes 0.50 and is not detected here. Closing that needs a scope INSIDE the callee") and [#1568](https://github.com/mudler/vllm.cpp/issues/1568) one level down. **MUST NOT be closed by a share floor whose denominator is the INSTRUMENT or the wall**: a seam's honest share is a property of the box, which makes that floor `residue <= 2 * instrument` with a different name, and `## Design` 3 of the spec is the measured record of why that does not work -- 4 red in 45 at the table bound, 28 in 160 on a standalone probe, and a 20-run sample that saw none of its tail. **The one shape that prohibition does NOT cover, and where a next attempt should start**: the anchor's own extent against the SEAM its two neighbours define, both measured in the same run -- near 1 honest, near 0 collapsed, with the region being named as the denominator instead of the instrument, so it never asks how many seconds a seam SHOULD hold. It still needs a constant and a cross-box distribution, so it is a candidate and not a conclusion. Otherwise: an anchor INSIDE the callee, or a bound on a quantity the scheduler cannot move ([#1570](https://github.com/mudler/vllm.cpp/issues/1570), [#1439](https://github.com/mudler/vllm.cpp/issues/1439)). NOT FIXED IN FLOW: each closure is its own row with its own red-first tests and fresh review. Listed under `## Owed` in [`ltx25-phase-residue.md`](../specs/ltx25-phase-residue.md) | bug | -| [#1906](https://github.com/mudler/vllm.cpp/issues/1906) | `LTX25-TEST-DETERMINISM` | **`test_ltx2_video`'s `/tmp` workspace is deleted by another process while a run is using it, and the abort costs 82 assertions with `failures="0"`.** The fixture writes every render into `"/tmp/vllm_ltx2_video_" + getpid() + "_" + counter` (`test_ltx2_video.cpp:67`), which is prefix-predictable and outside any per-run private directory. One run of seven at `ced0ab639` aborted with `cannot write /tmp/vllm_ltx2_video_970765_3/multichunk/audio.wav` after `create_directories` had succeeded and the frames had already been written to that same directory; it is not ENOSPC (49 GB free, and an ENOSPC open succeeds so the failure would read `short write`) and not a descriptor limit (`ulimit -n` 1048576, `file-nr` 6120). A sentinel planted at `/tmp/vllm_ltx2_video_999999_0` was GONE within two minutes while a second planted four minutes later survived twenty, so the deletion is episodic; `systemd-tmpfiles-clean` last fired 19 hours earlier, `grep -rn vllm_ltx2_video` over the tree returns only the fixture's own line, and a second agent was running its own build of this suite from another worktree throughout. An exception aborts a doctest case where it stands, so 89 of that case's 171 assertions ran and the run's total was 82 short with no failure recorded -- silent coverage loss wearing a normal number, and the second way this suite's total moves after [#1885](https://github.com/mudler/vllm.cpp/issues/1885). NOT fixed in flow: a random suffix under the same prefix still dies to `rm -rf /tmp/vllm_ltx2_video_*`, so where the fixture writes is a decision. Listed under `## Owed` in [ltx25-test-determinism.md](../specs/ltx25-test-determinism.md) | bug | -| [#1900](https://github.com/mudler/vllm.cpp/issues/1900) | `SPEC-DFLASH2` | **A non-causal SWA layer drops its window in our attention kernels; upstream attends within it.** Our sliding-window lower bound is guarded on `causal && window > 0` (`src/vt/cpu/cpu_ops.cpp:2917,2994` plus nine sites in `src/vt/cuda/cuda_ops.cu`, mirrored by `DflashBlockPagedMaskOf`), so a layer resolved `(causal=false, sliding_window=W)` attends the FULL context. At the pin, `vllm/model_executor/models/qwen3_dflash.py:89-146` resolves the window and the causal flag as two INDEPENDENT answers and `:221-234` passes `per_layer_sliding_window` irrespective of `causal`; this repository's own loader already says the same in prose (`src/vllm/model_executor/models/qwen3_dflash_weights.cpp:181-183`), so the kernels and that comment cannot both be right. Pre-existing and repo-wide, NOT introduced by #1890 — what W11 did was write it into a spec and a test header as a NORMATIVE claim without an upstream anchor, which is how it surfaced; both passages are corrected in [dflash2-draft-block-fa2.md](../specs/dflash2-draft-block-fa2.md) to describe the byte-identity the battery measures rather than to assert the semantics are right. Live because SPEC-DFLASH2 W1 (#1314) made `(causal=false, window>0)` reachable and `z-lab/Qwen3.8-27B-DFlash2` declares exactly that, so it is the production path for the published draft checkpoint: our draft attends over more context than upstream's, which moves acceptance and is invisible to a correctness gate because speculative decoding verifies every draft. NOT FIXED IN FLOW: it moves every attention kernel in the tree and needs its own red-before evidence. Listed under `## Owed` in [dflash2-draft-block-fa2.md](../specs/dflash2-draft-block-fa2.md) | bug | -| [#1902](https://github.com/mudler/vllm.cpp/issues/1902) | `SPEC-DFLASH2` | **W11's paged-seam guard never runs on a CUDA-graph REPLAY step, so the second refresh site is unguarded on every draft step after the first.** `detail::DflashBlockPagedAttention` re-derives the canonical `(slots, seq_ext)` pair from the store's own `ctx_len` and refuses a mismatch by name with no `kCPU` guard -- genuinely backend-independent, and it survives compiling out both `kCPU`-guarded reads beside it -- but the check is per CALL, not per STEP. On `st.g_state == 2` the driver calls `st.g_graph.Replay(queue)` and returns (`src/vllm/model_executor/models/qwen3_dflash.cpp:1636-1660`), so `ForwardPagedBody` and the guard with it are entered on the EAGER lane and on the ONE warm-then-capture step per request, and on NO replay step; those steps read persistent buffers refreshed by a SECOND production site (`:1627-1631`) with nothing downstream to check them. MEASURED by the fresh re-review of #1896: making that refresh skip on replay only (`if (st.g_state != 2) { ...Copy... }`), which on CUDA is a wrong answer at every step after the first, left ALL FOUR suites green (`decode_graph_seam` 4/4 23/23, `dflash2_draft` 43/43 449/449, `runner_reach` 8/8 162/162, `block_route` 13/13 30/30). Structurally invisible on CPU: the capture-capable CPU backend's `ReplayGraph` is a log push that executes nothing (`tests/vllm/models/decode_graph_seam_harness.h:117`), so the owed proof is DEVICE-SIDE. NOT FIXED IN FLOW: reading `g_seq_ext` back before `Replay` is a per-step D2H synchronisation on the path W11 exists to make faster, and moving the refresh inside the guarded function is impossible by construction because a replay never calls it and the buffer addresses are baked into the capture. The three places that overstated the guarantee are corrected. Listed under `## Owed` in [dflash2-draft-block-fa2.md](../specs/dflash2-draft-block-fa2.md) | bug | -| [#1901](https://github.com/mudler/vllm.cpp/issues/1901) | `MODEL-DSV4-EXL3` | **`exllamav3` is now pinned at `2398c05635fbbad01a0a51dce63c85c6c8a8450e` (tag `v1.4.3`) in [`.agents/oracles/exllamav3.md`](../oracles/exllamav3.md), and it records `gateable = no` because nothing here has ever BUILT or RUN it — this issue is what that field owes.** The registration itself closes a POLICY gap rather than a measurement one: every line of `MODEL-DSV4-EXL3` mirrors `turboderp-org/exllamav3` because vLLM registers no EXL3 at the parity pin, and AGENTS.md admits a secondary oracle "only when it appears in this table and has a recorded pin" — `exllamav3` appeared in neither the table nor `.agents/oracles/` while W1 landed on `main` and W2 sat in review. **No gate is invalidated today**: every W1/W2 gate runs against the in-tree CPU reference (`src/vt/cpu/cpu_exl3_dequant.cpp`) or an independently derived double-precision Sylvester H128, never against an exllamav3 execution. The row's own spec says so of its single real-checkpoint anchor — those spot values "were NOT produced by running upstream's kernel", because `ext.reconstruct` is a CUDA extension and the implementing host has no GPU — which removes transcription error from OUR side and cannot detect a defect in exllamav3 itself, since both sides read one source. **Why unmeasured**: `which nvcc` and `which nvidia-smi` both return nothing on the filing host, and `rc devices` read `dgx:gpu0 unhealthy (no contact 7m56s)` on 2026-08-25. **What would close it**: build exllamav3 at the pin for `sm_121a` in a `dgx:gpu0` lease, load `0xSero/deepseek-v4-flash-0731-spark`, and produce a greedy continuation. **A PREREQUISITE the oracle imposes, not a step**: `tp_import_split` is a runtime shared-memory IPC split (`exllamav3/model/model_tp_shared.py:23,40,285-292`; `exllamav3/modules/quant/exl3.py:285-329`) and NOTHING in its tree reads a `.rank{r}` filename, so the 43x4 rank-sliced artifact must be TP1-coalesced first — which is exactly what W1b's `LoadDeepseekV4Exl3` already produces — or the run needs four GPUs one Spark does not have. Both verdicts are results: if it runs the coalesced artifact the file flips to `gateable = yes` and it becomes the token oracle for W3c; if not, the file stays `no` with the reason recorded and the SparkInfer container run becomes the behavioral reference. **Two residues the same run closes**, both already under `## Owed` in [model-dsv4-exl3.md](../specs/model-dsv4-exl3.md): upstream's own `ext.reconstruct` against the W1a real-tensor anchors, and rank ORDER — nothing in this tree catches a transposed `.rank{r}` labelling, because four ranks concatenated in the wrong order reassemble to the right shape and pass every invariant, and no fixture can close it since the fixture writes the labels it then reads back. **One fact deliberately NOT asserted**: the checkpoint declares `hybrid_tr3_tail.exllamav3_revision = 787d1582…`, and whether that is an ancestor of the pin is UNVERIFIED here because the local clone is `depth = 1`, so its absence is a property of the clone and not evidence about upstream. NOT fixed in flow: it needs a GPU lease on a device that is currently unhealthy, a CUDA build of a third-party extension, and the ~107 GB checkpoint resident. Owned by row `MODEL-DSV4-EXL3` ([#1875](https://github.com/mudler/vllm.cpp/issues/1875)), where it falls due at **W3a** | bug | -| [#1922](https://github.com/mudler/vllm.cpp/issues/1922) | `ENG-POOL-BEST-FIT` | **`vllm::DevicePool` could only ever hand a freed scratch block back to a request in the block's OWN size class, so retention was a function of how many distinct shapes the traffic had shown rather than of how much one step concurrently needs — and nothing on the serving path ever gave any of it back.** Measured red-first on `main` @ `2e2b3fc1a`, CPU, through the production `LoadedEngine::generate` entry: twelve sequential requests with the LARGEST one FIRST, so every later request demanded strictly less than one already served and needed no buffer that request 0 had not already allocated and returned. The pool's retained bytes still went 560 243 -> 2 867 707 (**5.12x**) and it made **240 further driver allocations**, still climbing at request 11 (distinct classes 24 -> 101). With a DFlash2 draft attached — #1922's configuration in miniature — 622 120 -> 3 515 972 (5.65x). The three properties that make it unbounded were each already recorded and never joined up: blocks are never returned to the driver (the header says so), `Drain()` is called from exactly two places and neither is on the LLM path (`ltx2_video.cpp`, `minimax_h3_pipeline.cpp`), and `device_pool_cap_bytes` resolves to 0 — uncapped — on every platform. Because the per-step scratch shapes are `[T, hidden]`, `[C, kdim]`, `[S, vocab]` and so on, and how far `T` and `C` can travel is set by `max_model_len`, the ceiling rises with `--max-model-len`, which is #1922's own axis. **MIRROR:** torch's caching allocator, which is where vLLM's activations come from, searches its cache for the SMALLEST block at least as large as the request before it asks the driver (`c10/cuda/CUDACachingAllocator.cpp::get_free_block`); that search is the one structural difference, and its absence is what made a bounded upstream working set unbounded here. FIXED by giving `Get` that search over the class ladder, bounded at 2x so a caller never holds more than twice the bytes it asked for, and by recording in `block_class_` the class the driver actually allocated each live block at so `Put` returns a borrowed block to its OWN class — a loan, never a demotion, which is what stops a borrow starving the class it came from. Post-fix the same twelve requests retain 536 403 -> 624 899 (1.16x) with 27 driver allocations. A CUDA-graph capture cannot regress: `PreGrowForCapture` makes every class in the eager step's demand profile deep enough that every `Get` inside the capture HITS its own class, and a borrow is reachable only from a MISS. **THIS ROW DOES NOT CLOSE #1922**: the session had no GPU, so #1922's own `avail` curve at `--max-model-len 12288` on `dgx:gpu0` is neither reproduced nor attributed, and that run is `## Owed` O1. Spec [`pool-best-fit-retention.md`](../specs/pool-best-fit-retention.md) | bug | -| [#1926](https://github.com/mudler/vllm.cpp/issues/1926) | `ENG-POOL-BEST-FIT` | The other shape-keyed process-lifetime caches on the forward path, each keyed by an EXACT token count and each holding memory for every distinct count ever seen: `row_idx_by_t` (`dense_attn_block.h`, host, retained deliberately because a captured graph bakes its address), `DenseAlignFor` and its second copy in `qwen3_5.cpp` (four device buffers per distinct `M`, never freed), `MoeFusedResident`/`MoeBf16Resident` `tok_map`, and `cuda_matmul.cu`'s `heurs`/`plans` (whose `plans` values hold cuBLASLt descriptors that are never destroyed). Found while measuring #1922: with the `ENG-POOL-BEST-FIT` fix in — and equally with the pool switched off entirely under `VT_POOL_BYPASS=1`, the clean control — the same twelve requests still grow the heap by ~390 KB on a model whose hidden size is 32. An order of magnitude under the pool, and each needs its own decision (size class instead of exact count, a bound with eviction, or preallocation at `max_num_batched_tokens`); `row_idx_by_t` additionally needs the capture lifetime in the answer, so none of them is a one-line change. It is the floor under any memory-steady-state gate, which is why `ENG-POOL-BEST-FIT`'s gate reads `DevicePool::stats()` rather than process bytes. Listed under `## Owed` O2 in [`pool-best-fit-retention.md`](../specs/pool-best-fit-retention.md) | gap | -| [#1927](https://github.com/mudler/vllm.cpp/issues/1927) | `ENG-POOL-BEST-FIT` | `GPUModelRunner::connector_stored_blocks_` is a `request-id -> int` map that is inserted into and **never erased anywhere in the tree**, so a KV-connector server leaks one entry per request served for the life of the process. Found while auditing every request-keyed container in the runner for #1922. Honest scale: tens of bytes per request, and inert unless a worker-capable KV connector is installed, so it is NOT #1922's 2 GiB per request and is not claimed to be — but it is an unbounded request-keyed container in the request path, which is exactly the class of defect #1922 sent someone looking for, and the entry is dead the moment the request finishes (the count exists to deduplicate stores within ONE request's chunked prefill). NOT fixed in the `ENG-POOL-BEST-FIT` flow: it is a different behaviour with a different test surface — the gate is a KV-connector test that serves two requests and asserts the first is not retained — in a file that row's allocator change does not touch. Listed under `## Owed` O4 in [`pool-best-fit-retention.md`](../specs/pool-best-fit-retention.md) | bug | -| [#1510](https://github.com/mudler/vllm.cpp/issues/1510) | `LTX25-AUDIO-GUIDANCE-DEFAULTS` | **The LTX-2.5 guided audio arm is louder at the model's own defaults, and every step of the chain MIRRORS upstream, so this is upstream behavior and not a port defect.** Traced against Lightricks/LTX-2 `fd4ded7f2d88d3da713abcdd4ad41ecc4a9314ca` (registry id `ltx-2`, [`oracles/ltx-2.md`](../oracles/ltx-2.md), `gateable = no`) and corroborated at vLLM-Omni `a4ea67a21b20054dacc6e83952f9bd407e8ee4e7`: the four-term sum (`guiders.py:261-266` against `ltx2_pipeline.cpp::Ltx2MultiModalGuidance`), the standard-deviation renormalization (`guiders.py:268-271`), the audio guider defaults `cfg 7.0 / stg 1.0 / rescale 0.7 / modality 3.0` (`constants.py:59-68` against `ltx2_pipeline.cpp::Ltx2Params20`), combination in x0 space (`blocks.py:480-482`, `model.py:590-604`, the guider call at `denoisers.py:202-203`), the vocoder activation (`vocoder.py:436`) and the bandwidth-extension clamp with the int16 write (`vocoder.py:630`, `audio_mux.py:71-73`). Upstream has NO peak normalization anywhere (`normalize_audio_waveform` at `audio_mux.py:22-24` only transposes to channel-last) and NO clamp on the latent or on the guided prediction. The renormalization hypothesis in the issue is FALSE: `rescale_scale` is 0.7 on BOTH arms. vLLM-Omni's `_official_guidance` (`ltx2_recipes.py:90-106`) ships the same eight values. Upstream DOCUMENTS the trade: `multimodal-guidance.md:14` calls `rescale_scale` the knob that "Helps prevent over-saturation" with a typical band of 0.5-0.7, and `:11` gives `cfg_scale` a typical band of 2.0-5.0 while `constants.py:61` ships the AUDIO arm at 7.0. RE-MEASURED: the issue's unguided envelope CV of 0.1069 is a transposition of the AAC value 0.109629 (`verify.json` records 0.11184659763714917), the compression mechanism it asserts is absent (2 samples of 96,960 at full scale, longest run 1 sample, whole distribution scaled ~1.45x), and the A/B is confounded by a different binary and a `--checkpoint-class full` flag beside the one under test. THE REAL GAP, and what this row builds: the AUDIO guider's resolved scales had NO render-level assertion while the video guider's four did, so disabling the audio renormalization left the one_stage X0-space case GREEN at 67/67. Adds four `audio_guidance_*` fields to `Ltx2ConditioningTrace` and pins them against the upstream literals. NO render behavior changes. Spec [`ltx25-audio-guidance-defaults.md`](../specs/ltx25-audio-guidance-defaults.md) | bug | -| [#1905](https://github.com/mudler/vllm.cpp/issues/1905) | `LTX25-AUDIO-GUIDANCE-DEFAULTS` | **`verify_render.py` produced the FAIL verdict [#1510](https://github.com/mudler/vllm.cpp/issues/1510) rests on, and it is not in this repository.** Its only copy is on a CIFS share, `/mnt/nas_share/rc/ltx25-fullmodel/job/verify_render.py`, sha256 `57cf92846506be961e3c6ab3c9198de5608d0e85bfd1eb992eb91dda1c0fa563`. It is unversioned, untested and unreachable from any gate here; `grep -rn 'envelope_cv'` over this tree returns nothing. Two properties bear on the verdicts it produced. Its `envelope_cv` rests on 20 non-overlapping 50 ms frames over a 1.01 s clip (`hop = max(int(sr * 0.05), 1)` at 48 kHz over 48,480 frames) against a FAIL threshold of 0.10, with no confidence interval computed or reported. Its `active_fraction` threshold is RELATIVE, at `env.max() * 10 ** (-40 / 20.0)`, so a constant-level signal reads 1.0 at any absolute level, which is why both arms of the #1510 A/B read exactly 1.0 while the issue table shows a dash against 1.0. It also downmixes to mono before every headline metric, so its `peak_dbfs` of -0.287 is a mixdown value against a stereo peak sample of 32767. Until the instrument is in the tree with a test and a gate, no audio PASS or FAIL from this campaign is a gate result. Listed under `## Owed` in [`ltx25-audio-guidance-defaults.md`](../specs/ltx25-audio-guidance-defaults.md) | bug | -| [#1920](https://github.com/mudler/vllm.cpp/issues/1920) | `LTX25-AUDIO-GUIDANCE-DEFAULTS` | **The audio guider's `skip_step` is the one field of `denoise_in.audio_guider` that nothing observes.** [#1510](https://github.com/mudler/vllm.cpp/issues/1510) closed the CONSUMPTION gap by recording the four `audio_first_*` passes and replaying `ltx2_pipeline.cpp::Ltx2MultiModalGuidance` over them, exactly, at the recipe's own scales and again at four request overrides. That replay bites -- handing the denoiser a copy with `cfg_scale = 1.0, rescale_scale = 0.0` reds both cases -- but it holds only what the function READS. `Ltx2MultiModalGuiderParams` carries six fields: `cfg_scale`, `stg_scale`, `rescale_scale` and `modality_scale` are held by the replay; `stg_blocks` is held separately by `Ltx2ConditioningTrace::video_audio_perturbed_blocks`, which the denoiser derives from the same params (`ltx2_denoisers.cpp::Ltx2GuidedDenoise`); `skip_step` is held by **nothing**. `ltx2_pipeline.cpp::ShouldSkipStep` is `step % (skip_step + 1) != 0` (`guiders.py:287-291`), which is FALSE at step 0 for every `skip_step`, and `Ltx2ConditioningTrace` describes the FIRST guided step only -- so an `audio_guider` mis-handed on that field alone moves no recorded tensor and every assertion stays green. The shipped defaults do not close it either: `grep -n 'skip_step' src/vllm/model_executor/models/ltx2_pipeline.cpp` finds four assignments, all `= 0`, so a cross-wire between the video and audio guiders is invisible on a default render for the same reason `stg_scale`, `rescale_scale` and `modality_scale` were before #1510's override case existed: the wrong source carries the right number. **What would close it**, two pieces and neither is a comment: a trace field for the skip decision -- `Ltx2GuidedDenoiseResult` already carries `audio_skipped` and `ltx2_video.cpp::RecordFirstGuidedStep` copies it nowhere -- recording a step the guider CAN skip, since step 0 never is; and a render that reaches one, which the existing `kLtx2AudioSkipStepExtra` request extra makes reachable at `audio_skip_step = 1` on the two-step fixture renders. Its own red-first mutation is named: hand `denoise_in.audio_guider` a copy whose `skip_step` is the video row's while every scale stays correct. Found by the fresh review of #1510's repair. NOT FIXED IN FLOW: it adds a field to a shared struct, needs a multi-step assertion the trace cannot express today, and needs its own red-before evidence. No render behavior changes and no gate that exists today is invalidated. Listed under `## Owed` in [`ltx25-audio-guidance-defaults.md`](../specs/ltx25-audio-guidance-defaults.md) | bug | -| [#1919](https://github.com/mudler/vllm.cpp/issues/1919) | `SPEC-DFLASH2` | **The DFlash2 draft context store is hard-capped at 4096 slots, so a >4K-token prompt kills EngineCore and every later request on that server gets `[request submitted to a stopped AsyncLLM]`.** `kDflashMaxCtxSlots = 4096` (`src/vllm/model_executor/models/qwen3_dflash.cpp:1006`) sizes `max_pages` regardless of `--max-model-len`, so the engine advertises 12288, admits the request, and then throws `AppendContextKVDevice: paged store capacity exceeded` from inside the EngineCore step. Upstream has no private store and no private cap: the DFlash draft's context K/V goes into the engine's own paged KV cache (`vllm/model_executor/models/qwen3_dflash.py:604-620` at pin `5559679229`), whose block tables are `cdiv(max_model_len, block_size)` (`vllm/v1/worker/gpu/model_runner.py:426,444`); and where a speculator cannot serve a request it emits an EMPTY draft and lets the target run alone (`vllm/v1/spec_decode/ngram_proposer.py:156-159`, `suffix_decoding.py:59-62`), never raising. Repair: size the store from `max_model_len + num_query_per_req` under a per-request byte budget, fall back to the non-speculative path for a request that outgrows it, and announce the effective speculative context once at startup — wave spec [dflash2-ctx-store-capacity.md](../specs/dflash2-ctx-store-capacity.md) | bug | -| [#1943](https://github.com/mudler/vllm.cpp/issues/1943) | `SPEC-DFLASH2` | **A DFlash2 request that has fallen back under async scheduling keeps paying a full `1 + k` verify at ~zero acceptance for the rest of its life; upstream TRIMS the schedule instead.** [#1919](https://github.com/mudler/vllm.cpp/issues/1919) made a request whose context outgrows the draft store fall back rather than kill EngineCore, and under async scheduling the empty draft is not available: `AsyncScheduler::update_after_schedule` places `num_spec_tokens_to_schedule` placeholders one step BEFORE the propose that fills them and `Scheduler::update_draft_token_ids` is never called under async (`src/vllm/v1/engine/core.cpp:120-123`), so the runner keeps the draft's SHAPE and neutralises its CONTENT with the draft's mask token. That is CORRECT — the verify is lossless, so the emitted stream is identical and a token gate cannot see this at all — and it is WASTEFUL: the request is scheduled `1 + k` verify positions on every later step and accepts essentially none of them, about 9x the target compute per emitted token at `k = 8`. Upstream's DFlash draft keeps no private store (`vllm/model_executor/models/qwen3_dflash.py:604-620` at pin `5559679229`) so it has no fallback state to mirror; its nearest mechanism for a proposer that delivered fewer drafts than the scheduler budgeted is `update_scheduler_for_invalid_drafts` (`vllm/v1/spec_decode/ngram_proposer_gpu.py:475-515`, called from `vllm/v1/worker/gpu_model_runner.py:1333-1344`), which decrements `num_scheduled_tokens` and `total_num_scheduled_tokens` and pops the request out of `scheduled_spec_decode_tokens` at `valid_k == 0`, keeping `original_num_spec_per_req` for the rejection correction. NOT FIXED IN FLOW: upstream gates that trim on `use_ngram_gpu()` and applies it to neither the eagle nor the DFlash family, so porting it here generalises an upstream mechanism rather than transcribing upstream's DFlash arm; it moves the scheduler/worker contract; and it needs its own red-before gate on the SCHEDULE. The alternative already owed — moving the draft's context K/V into the engine's own paged allocator — subsumes it. Found by the fresh review of [#1932](https://github.com/mudler/vllm.cpp/pull/1932). Listed under `## Owed` in [dflash2-ctx-store-capacity.md](../specs/dflash2-ctx-store-capacity.md) | bug | +| [#1434](https://github.com/mudler/vllm.cpp/issues/1434) | `GATE-DOC-CHECKPOINT-STATES` | **`scripts/check-doc-checkpoint.py` could not see `PARTIAL`, so 118 state cells could move with no gate observing them.** `STATES` (`:56-66`) is the whole definition of what a lifecycle state IS for the gate that enforces AGENTS.md's `docs/STATUS.md` / `docs/BENCHMARKS.md` / spec `## Now` triple, and `row_states` drops any row it cannot match. `lifecycle_moves` and `moved_rows` then iterate the AFTER map, so leaving the matched set is silent by construction. Re-derived at `947e5f648` (the report measured `63d87805c`): `PARTIAL` **118** cells and `ANCHOR-BACKFILL` **73**, against `DONE` 77 and `BLOCKED` 9 — `PARTIAL` is the second most used state in the matrices and the gate was blind to it. Over the seven tables `ROW_TABLES` actually reads, the resolved population goes from **153 rows to 226**, a 47.7 % widening. Two of the transitions the report names behave differently from its description, measured with scratch commits at `947e5f648` on an unmodified checker: `READY -> PARTIAL` rc **0** and `PARTIAL -> READY` rc **0** are the real blind spots, while the report's suggested `PARTIAL -> ACTIVE` already reds — by accident, reporting **`added as ACTIVE`** for a row that has existed for months, because it is absent from the BEFORE map. FIXED IN FLOW for `PARTIAL` only. **`ANCHOR-BACKFILL` is deliberately excluded**: `.agents/feature-matrix.md:14-17` defines it as a property of the RECORD (*a legacy implemented row without exact code, test and real-spec anchors*), `docs/STATUS.md` carries no such term and would have nothing true to write on a `DONE <-> ANCHOR-BACKFILL` move, and `REQUIRED["lifecycle"]` cannot demand the spec's `## Now` alone — so admitting it would demand a public-document edit with nothing to say, which is the exact shape `check-doc-checkpoint.py:4-17` records as the reason the file was rewritten (16 of 20 red CI runs, six hardcoded escape hatches). One row's resolved state moves and the move is a REPAIR: `KV-BLOCK-POOL` says `` `PARTIAL` (not `DONE`) `` in its prose and the last-match heuristic believed the parenthesis, resolving `DONE`. No pinned counter moves — `check-gate-commands.py` has its own `GATED_STATES` and `RUNNABLE_BASELINE` is keyed on matrix rows, `UNOWNED_HIGH_WATER` is unmoved because this row names an owner, and no matrix row or public document changes — which was measured, not assumed, because this is the [#1376](https://github.com/mudler/vllm.cpp/issues/1376) ratchet shape. Remainder listed under `## Owed` in [doc-checkpoint-lifecycle-states.md](specs/doc-checkpoint-lifecycle-states.md): `ANCHOR-BACKFILL` moves, `.agents/sglang-matrix.md` never entering `ROW_TABLES`, a row that leaves the matched set entirely, and a new row added directly as `PARTIAL` | bug | +| [#1586](https://github.com/mudler/vllm.cpp/issues/1586) | `BACKEND-ROCM` | Adopt the ROCm 7.14 container toolchain (first TheRock production release, gfx1100 supported) and open the gfx1100 optimization campaign; baseline recorded in the issue: build 586/586 green on `rocm-dev:7.14.0`, focused gate 4/5 with the `MoeSiluMul` bf16 exactness failure at `tests/vt/test_backend_cross_device.cpp:2063` | feature | +| [#1587](https://github.com/mudler/vllm.cpp/issues/1587) | `KERNEL-QUANT-CIQ-GEMM-ROCM` | The ROCm backend has no quantized-weight GEMM provider, so every GGUF k-quant on an AMD card computes off device; upstream pins a `csrc/rocm` W4A16 family (`gptq_gemm_rdna3` gated `VLLM_ROCM_GFX1100`) we cannot reach yet because no GPTQ consumer exists. W1 registers the `kMatmulBTQuant`/`kMatmulBTQuantGrouped` keep-quant providers mirroring the CUDA sibling; W2 (owed) ports the upstream family behind a loader consumer. Spec [`kernel-quant-ciq-gemm-rocm.md`](specs/kernel-quant-ciq-gemm-rocm.md) | feature | +| [#1588](https://github.com/mudler/vllm.cpp/issues/1588) | `BACKEND-ROCM` | Characterize Qwen3.5-0.8B CPU against ROCm numerics on gfx1100: the backend matrix records the all-native run with its numerical characterization open; owns the `MoeSiluMul` bf16 exactness failure found in the 7.14 baseline | verification | +| [#7](https://github.com/ghazni101/vllm.cpp/issues/7) | `GFX1100-TG200` | ROCm gfx1100: fp8 KV cache disables all optimized decode attention kernels — `PagedAttnOnline` fallback is 7.5x slower than `PagedAttnDecodeGqaF32Q` at 16K context. Widen the `VT_ATTN_DECODE_GQA4` dtype guard from `kBF16` to accept `kI8` KV as well and add an fp8 dequant load path in the GQA kernel. Spec PENDING on this branch; the W6 fp8 decode-attn extras are parked on branch `row/fp8-kv-decode-attn` (`3a345b5ae`, 2026-08-29) | perf | +| [#1919](https://github.com/mudler/vllm.cpp/issues/1919) | `SPEC-DFLASH2` | **The DFlash2 draft context store is hard-capped at 4096 slots, so a >4K-token prompt kills EngineCore and every later request on that server gets `[request submitted to a stopped AsyncLLM]`.** `kDflashMaxCtxSlots = 4096` (`src/vllm/model_executor/models/qwen3_dflash.cpp:1006`) sizes `max_pages` regardless of `--max-model-len`, so the engine advertises 12288, admits the request, and then throws `AppendContextKVDevice: paged store capacity exceeded` from inside the EngineCore step. Upstream has no private store and no private cap: the DFlash draft's context K/V goes into the engine's own paged KV cache (`vllm/model_executor/models/qwen3_dflash.py:604-620` at pin `5559679229`), whose block tables are `cdiv(max_model_len, block_size)` (`vllm/v1/worker/gpu/model_runner.py:426,444`); and where a speculator cannot serve a request it emits an EMPTY draft and lets the target run alone (`vllm/v1/spec_decode/ngram_proposer.py:156-159`, `suffix_decoding.py:59-62`), never raising. Repair: size the store from `max_model_len + num_query_per_req` under a per-request byte budget, fall back to the non-speculative path for a request that outgrows it, and announce the effective speculative context once at startup — wave spec [dflash2-ctx-store-capacity.md](specs/dflash2-ctx-store-capacity.md) | bug | +| [#1943](https://github.com/mudler/vllm.cpp/issues/1943) | `SPEC-DFLASH2` | **A DFlash2 request that has fallen back under async scheduling keeps paying a full `1 + k` verify at ~zero acceptance for the rest of its life; upstream TRIMS the schedule instead.** [#1919](https://github.com/mudler/vllm.cpp/issues/1919) made a request whose context outgrows the draft store fall back rather than kill EngineCore, and under async scheduling the empty draft is not available: `AsyncScheduler::update_after_schedule` places `num_spec_tokens_to_schedule` placeholders one step BEFORE the propose that fills them and `Scheduler::update_draft_token_ids` is never called under async (`src/vllm/v1/engine/core.cpp:120-123`), so the runner keeps the draft's SHAPE and neutralises its CONTENT with the draft's mask token. That is CORRECT — the verify is lossless, so the emitted stream is identical and a token gate cannot see this at all — and it is WASTEFUL: the request is scheduled `1 + k` verify positions on every later step and accepts essentially none of them, about 9x the target compute per emitted token at `k = 8`. Upstream's DFlash draft keeps no private store (`vllm/model_executor/models/qwen3_dflash.py:604-620` at pin `5559679229`) so it has no fallback state to mirror; its nearest mechanism for a proposer that delivered fewer drafts than the scheduler budgeted is `update_scheduler_for_invalid_drafts` (`vllm/v1/spec_decode/ngram_proposer_gpu.py:475-515`, called from `vllm/v1/worker/gpu_model_runner.py:1333-1344`), which decrements `num_scheduled_tokens` and `total_num_scheduled_tokens` and pops the request out of `scheduled_spec_decode_tokens` at `valid_k == 0`, keeping `original_num_spec_per_req` for the rejection correction. NOT FIXED IN FLOW: upstream gates that trim on `use_ngram_gpu()` and applies it to neither the eagle nor the DFlash family, so porting it here generalises an upstream mechanism rather than transcribing upstream's DFlash arm; it moves the scheduler/worker contract; and it needs its own red-before gate on the SCHEDULE. The alternative already owed — moving the draft's context K/V into the engine's own paged allocator — subsumes it. Found by the fresh review of [#1932](https://github.com/mudler/vllm.cpp/pull/1932). Listed under `## Owed` in [dflash2-ctx-store-capacity.md](specs/dflash2-ctx-store-capacity.md) | bug | | [#1867](https://github.com/mudler/vllm.cpp/issues/1867) | `SPEC-DFLASH2` | **The DFlash2 selector's per-step top-k cost 683 us/step where FlashInfer's radix top-k does the same work in 40 us** -- `TopKValuesIndicesRowKernel`, 8 rows x 248320 vocab, K=16, measured on `dgx:gpu0` with nsys against SGLang on the identical checkpoint and workload (#1857's kernel table, the artifact-verified re-take): +0.65 ms/step, the fourth-largest per-step lever there. The cost was the ITERATION COUNT -- a ternary bisection of the threshold in float VALUE space under `kThreshMaxIter = 64`, every iteration a full pass over a 248320-wide row -- where a radix narrowing over a monotone key fixes the same threshold EXACTLY in four rounds, two of which read global memory at all. W12 ports the arithmetic (`include/vt/radix_topk.h`, anchored on `flashinfer/topk_common.cuh:35-39` and `flashinfer/topk.cuh:683-691` at FlashInfer `0.6.12`, the wheel vLLM's own `_topk` dispatches to at merge `b389ac29`) and rewrites the CUDA arm around it as `TopKValuesIndicesRadixRowKernel`; the multi-CTA grid barrier and workspace `## Risks/decisions` D2 refused stay refused, ONE CTA PER ROW. **The tie-break does not move**: upstream leaves FlashInfer's `tie_break` at `NONE`, ours is index-ascending and `include/vt/ops.h` pins it, so the port mirrors FlashInfer's algorithm and our contract -- which is FlashInfer's own `TopKTieBreak::Small`. The CPU reference is UNCHANGED, so the two arms still answer by different routes. Gated on a host with no `nvcc` by `tests/vt/test_ops_radix_topk` against a full stable sort, including on the production shape. **The GPU number and the device run are OWED** (`## Owed` O34, operator-run, `-DVLLM_CPP_CUTLASS_FETCH=ON` plus an `nm` assertion on the new kernel name before any timing); occupancy is the named residual (O35). Nothing here claims a measured speedup | perf | -| [#1923](https://github.com/mudler/vllm.cpp/issues/1923) | `MODEL-DSV4-EXL3` | **The EXL3 loader never set `has_host_weights`, so a loaded EXL3 tower could never be consumed — and the W2 forward tests hid it by setting the flag by hand.** `LoadDeepseekV4Exl3` accounted for the `carried-*` tensors with `require(...)` -> `RequireTensor`, a presence check that increments a counter; it wrote nothing into `DeepseekV4Weights::host` and never set the flag every forward entry point gates on, so `has_exl3_weights && has_host_weights` was unreachable from any load and an end-to-end `vllm-server` probe generated ZERO tokens. Three review rounds and a mutation pass missed it because `tests/vllm/models/test_deepseek_v4_exl3_forward.cpp` set `has_host_weights = true` itself at five sites and built `DeepseekV4Weights` directly instead of going through `LoadDeepseekV4ForCausalLMWeights` — `.agents/reachability.md`'s documented failure in its exact shape. **FIXED IN FLOW (W1c).** The loader now MATERIALIZES the carried half into the host-float tower `ForwardComposeImpl` composes with: block-wise FP8 (`F8_E4M3` weight + `F8_E8M0` scale over the artifact's own `weight_block_size` `[128,128]`) decoded to f32 through a new shared `DequantFp8BlockToF32` in the fp8/nvfp4 dequant family, BF16 norms/embeddings/router widened, F32 MHC and sinks straight through, I64 `tid2eid` narrowed to int32 — each destination shape derived from the resolved config and REFUSED BY NAME on a mismatch, because `Gemm`'s host arm is a `MatVec` with no length check and a wrong shape there is a silently wrong number. Both suites now drive the PRODUCTION loader over ONE shared hermetic fixture written at the real dtypes (`tests/vllm/models/dsv4_exl3_fixture.h`), and a synthetic rank-sliced checkpoint LOADS and EMITS LOGITS end to end (`rel_rms` 1.68e-3 vs the dequantized-dense reference, 4.32 vs an unrelated one). The issue's second finding is settled by DELETION: the EXL3-specific row-naming refusal at `deepseek_v4.cpp` was unreachable on the default path (the runner's default `gather` routes to `ForwardDevice`, whose generic `kHostPending` fires first) and W1c makes the state it guarded unreachable from ANY load, since the one arm that sets `has_exl3_weights` now sets `has_host_weights` before returning; `MoeBlock` gains the guard that actually pays, refusing by name when the host routed-expert tower is empty, which is what turns the reachability mutation's verdict from freed-memory UB into a message. The residency refusal now prices the carried tower too (~29 GB of f32 beside ~84 GiB of trellis on the real artifact). **STILL OPEN, and this issue carries it**: the REAL artifact does not run. Its DSA tensors are stored at twice the width the host forward indexes (`compressor.wgate` `[2*head_dim, H]`, `indexer.compressor.wkv` `[2*index_head_dim, H]`, `indexer.wq_b` `[inh*ihd, q_lora_rank]`, all MEASURED 2026-08-25) and 41 of its 43 layers carry a compressor, so the loader refuses. The GGUF arm dodges the identical geometry with `dsa_dense = (be.gguf != nullptr)`; widening that predicate to the EXL3 source would break this row's own equivalence gate, whose dense reference has `be.exl3 == nullptr` and would stop taking the same attention path. The fix is a dense-MLA selector BOTH arms read, which is a DSA-residual policy decision no row owns. The real artifact is additionally blocked by its tokenizer ([#1924](https://github.com/mudler/vllm.cpp/issues/1924)). That residual is owned by row `MODEL-DSV4-EXL3` and listed under `## Owed` in [model-dsv4-exl3.md](../specs/model-dsv4-exl3.md), NOT by this issue, which is the loader defect and is closed by W1c | bug | -| [#1925](https://github.com/mudler/vllm.cpp/issues/1925) | `KV-DSV4-MULTICACHE` | **`DeepseekV4Model::Forward` and `ForwardDevice` both open `(void)attn_meta; (void)attn_kv;` (`src/vllm/model_executor/models/deepseek_v4.cpp:2886-2887`, `:2959-2960`), so the paged KV cache the runner hands in is discarded and every server-side decode step recomputes the whole prefix.** `MakeDeepseekV4KVCache` (`deepseek_v4_registry.cpp:126-148`) says so itself -- "STUB (W3) ... a multi-cache geometry not yet representable ... Never exercised this pass". SCOPED 2026-08-25 by [kv-dsv4-multicache.md](../specs/kv-dsv4-multicache.md), which derives the geometry from both sides at the pin `5559679229bc961848b121ccdeaa8fa5d79bec98`: DeepSeek-V4-Flash needs **167 cache entries across 43 layers** -- 43 `SlidingWindowMLASpec` SWA caches (block_size **64**, window 128, on every attention layer including the two with `compress_ratio == 0`, which have no MLA cache at all), 41 `MLAAttentionSpec` compressed latents at **584 B/token** (448B NoPE + 8B UE8M0 scale + 128B RoPE), 21 indexer key caches (132 B FP8 / 68 B MXFP4), and 62 f32 compressor state caches at block_size **4 or 8** (41 on the attention compressors, 21 more on the indexers' own). **"Not representable" is all four failures at once, not one:** our `MLAAttentionSpec` (`include/vllm/v1/kv_cache_interface.h:242-261`) adds NO fields over `FullAttentionSpec` and so carries neither `compress_ratio` (which upstream makes `storage_block_size = block_size // compress_ratio`, `kv_cache_interface.py:394-395`) nor `alignment`; `SlidingWindowMLASpec` does not exist at all -- the enumerator `kSlidingWindowMla` is declared with no struct behind it and the port's own deferral list names the class as omitted (`kv_cache_interface.h:46-52`), which is 105 of the 167 entries; a `KVCacheGroupSpec` holds exactly one spec (`:358-369`); and the runner keeps at most ONE attention group and ONE recurrent group, dropping every other kind **with no diagnostic** (`src/vllm/v1/worker/gpu/runner.cpp:577-597`, three-valued `LayerKvClass` at `include/vllm/v1/worker/gpu/runner.h:366-370`). `ModelForwardInput` also has no third channel to receive them (`include/vllm/model_executor/models/model_registry.h:303-315`), and non-uniform `block_size` across groups is deferred by an `assert` -- inert under `NDEBUG` -- at `src/vllm/v1/core/kv_cache_coordinator.cpp:340-346`. **The consequence for measurement, which is why it was asked now:** a server tok/s figure here is not a slow decode rate, it is a different quantity -- per-token cost grows with position. The one arm that does cache, `DeepseekV4ForwardGgufCached`, runs `dsa_dense` with the indexer and compressor forced OFF on every layer (`deepseek_v4.cpp:677-679`) and refuses otherwise (`VT_CHECK(!is_indexer && !is_comp, ...)`, `:786-787`), which the code's own comment (`:664-676`) says is exact only while `seq_len <= index_topk` (=512); above that it attends over a token set upstream does not. **A token gate at or below 512 tokens therefore cannot detect this defect.** TWO FINDINGS THE SCOPING ADDS beyond the issue text: the attention sink is a per-head weight loaded from the checkpoint (`attention.py:218-222`, ours `include/vllm/model_executor/models/deepseek_v4.h:159`) and NOT cache state, so it is out of scope; and Gemma-3's ratified "sliding window is a mask, not a cache" reasoning (`src/vllm/model_executor/models/gemma3_registry.cpp:105-109`) does NOT carry to V4, whose SWA cache is the only cache on layers 0 and 1 and whose 64-token page is fixed by physical tensor sharing with the C4A blocks (`sparse_swa.py:76-83`). NOT FIXED IN FLOW and deliberately: the deliverable asked for was a scoping document, and the work is a multi-wave capability across the KV interface, the runner and the model, whose W3 touches every model's allocation path. Owned by row `KV-DSV4-MULTICACHE` and listed under `## Owed` in [kv-dsv4-multicache.md](../specs/kv-dsv4-multicache.md) | bug | +| [#1923](https://github.com/mudler/vllm.cpp/issues/1923) | `MODEL-DSV4-EXL3` | **The EXL3 loader never set `has_host_weights`, so a loaded EXL3 tower could never be consumed — and the W2 forward tests hid it by setting the flag by hand.** `LoadDeepseekV4Exl3` accounted for the `carried-*` tensors with `require(...)` -> `RequireTensor`, a presence check that increments a counter; it wrote nothing into `DeepseekV4Weights::host` and never set the flag every forward entry point gates on, so `has_exl3_weights && has_host_weights` was unreachable from any load and an end-to-end `vllm-server` probe generated ZERO tokens. Three review rounds and a mutation pass missed it because `tests/vllm/models/test_deepseek_v4_exl3_forward.cpp` set `has_host_weights = true` itself at five sites and built `DeepseekV4Weights` directly instead of going through `LoadDeepseekV4ForCausalLMWeights` — `.agents/reachability.md`'s documented failure in its exact shape. **FIXED IN FLOW (W1c).** The loader now MATERIALIZES the carried half into the host-float tower `ForwardComposeImpl` composes with: block-wise FP8 (`F8_E4M3` weight + `F8_E8M0` scale over the artifact's own `weight_block_size` `[128,128]`) decoded to f32 through a new shared `DequantFp8BlockToF32` in the fp8/nvfp4 dequant family, BF16 norms/embeddings/router widened, F32 MHC and sinks straight through, I64 `tid2eid` narrowed to int32 — each destination shape derived from the resolved config and REFUSED BY NAME on a mismatch, because `Gemm`'s host arm is a `MatVec` with no length check and a wrong shape there is a silently wrong number. Both suites now drive the PRODUCTION loader over ONE shared hermetic fixture written at the real dtypes (`tests/vllm/models/dsv4_exl3_fixture.h`), and a synthetic rank-sliced checkpoint LOADS and EMITS LOGITS end to end (`rel_rms` 1.68e-3 vs the dequantized-dense reference, 4.32 vs an unrelated one). The issue's second finding is settled by DELETION: the EXL3-specific row-naming refusal at `deepseek_v4.cpp` was unreachable on the default path (the runner's default `gather` routes to `ForwardDevice`, whose generic `kHostPending` fires first) and W1c makes the state it guarded unreachable from ANY load, since the one arm that sets `has_exl3_weights` now sets `has_host_weights` before returning; `MoeBlock` gains the guard that actually pays, refusing by name when the host routed-expert tower is empty, which is what turns the reachability mutation's verdict from freed-memory UB into a message. The residency refusal now prices the carried tower too (~29 GB of f32 beside ~84 GiB of trellis on the real artifact). **STILL OPEN, and this issue carries it**: the REAL artifact does not run. Its DSA tensors are stored at twice the width the host forward indexes (`compressor.wgate` `[2*head_dim, H]`, `indexer.compressor.wkv` `[2*index_head_dim, H]`, `indexer.wq_b` `[inh*ihd, q_lora_rank]`, all MEASURED 2026-08-25) and 41 of its 43 layers carry a compressor, so the loader refuses. The GGUF arm dodges the identical geometry with `dsa_dense = (be.gguf != nullptr)`; widening that predicate to the EXL3 source would break this row's own equivalence gate, whose dense reference has `be.exl3 == nullptr` and would stop taking the same attention path. The fix is a dense-MLA selector BOTH arms read, which is a DSA-residual policy decision no row owns. The real artifact is additionally blocked by its tokenizer ([#1924](https://github.com/mudler/vllm.cpp/issues/1924)). That residual is owned by row `MODEL-DSV4-EXL3` and listed under `## Owed` in [model-dsv4-exl3.md](specs/model-dsv4-exl3.md), NOT by this issue, which is the loader defect and is closed by W1c | bug | +| [#1925](https://github.com/mudler/vllm.cpp/issues/1925) | `KV-DSV4-MULTICACHE` | **`DeepseekV4Model::Forward` and `ForwardDevice` both open `(void)attn_meta; (void)attn_kv;` (`src/vllm/model_executor/models/deepseek_v4.cpp:2886-2887`, `:2959-2960`), so the paged KV cache the runner hands in is discarded and every server-side decode step recomputes the whole prefix.** `MakeDeepseekV4KVCache` (`deepseek_v4_registry.cpp:126-148`) says so itself -- "STUB (W3) ... a multi-cache geometry not yet representable ... Never exercised this pass". SCOPED 2026-08-25 by [kv-dsv4-multicache.md](specs/kv-dsv4-multicache.md), which derives the geometry from both sides at the pin `5559679229bc961848b121ccdeaa8fa5d79bec98`: DeepSeek-V4-Flash needs **167 cache entries across 43 layers** -- 43 `SlidingWindowMLASpec` SWA caches (block_size **64**, window 128, on every attention layer including the two with `compress_ratio == 0`, which have no MLA cache at all), 41 `MLAAttentionSpec` compressed latents at **584 B/token** (448B NoPE + 8B UE8M0 scale + 128B RoPE), 21 indexer key caches (132 B FP8 / 68 B MXFP4), and 62 f32 compressor state caches at block_size **4 or 8** (41 on the attention compressors, 21 more on the indexers' own). **"Not representable" is all four failures at once, not one:** our `MLAAttentionSpec` (`include/vllm/v1/kv_cache_interface.h:242-261`) adds NO fields over `FullAttentionSpec` and so carries neither `compress_ratio` (which upstream makes `storage_block_size = block_size // compress_ratio`, `kv_cache_interface.py:394-395`) nor `alignment`; `SlidingWindowMLASpec` does not exist at all -- the enumerator `kSlidingWindowMla` is declared with no struct behind it and the port's own deferral list names the class as omitted (`kv_cache_interface.h:46-52`), which is 105 of the 167 entries; a `KVCacheGroupSpec` holds exactly one spec (`:358-369`); and the runner keeps at most ONE attention group and ONE recurrent group, dropping every other kind **with no diagnostic** (`src/vllm/v1/worker/gpu/runner.cpp:577-597`, three-valued `LayerKvClass` at `include/vllm/v1/worker/gpu/runner.h:366-370`). `ModelForwardInput` also has no third channel to receive them (`include/vllm/model_executor/models/model_registry.h:303-315`), and non-uniform `block_size` across groups is deferred by an `assert` -- inert under `NDEBUG` -- at `src/vllm/v1/core/kv_cache_coordinator.cpp:340-346`. **The consequence for measurement, which is why it was asked now:** a server tok/s figure here is not a slow decode rate, it is a different quantity -- per-token cost grows with position. The one arm that does cache, `DeepseekV4ForwardGgufCached`, runs `dsa_dense` with the indexer and compressor forced OFF on every layer (`deepseek_v4.cpp:677-679`) and refuses otherwise (`VT_CHECK(!is_indexer && !is_comp, ...)`, `:786-787`), which the code's own comment (`:664-676`) says is exact only while `seq_len <= index_topk` (=512); above that it attends over a token set upstream does not. **A token gate at or below 512 tokens therefore cannot detect this defect.** TWO FINDINGS THE SCOPING ADDS beyond the issue text: the attention sink is a per-head weight loaded from the checkpoint (`attention.py:218-222`, ours `include/vllm/model_executor/models/deepseek_v4.h:159`) and NOT cache state, so it is out of scope; and Gemma-3's ratified "sliding window is a mask, not a cache" reasoning (`src/vllm/model_executor/models/gemma3_registry.cpp:105-109`) does NOT carry to V4, whose SWA cache is the only cache on layers 0 and 1 and whose 64-token page is fixed by physical tensor sharing with the C4A blocks (`sparse_swa.py:76-83`). NOT FIXED IN FLOW and deliberately: the deliverable asked for was a scoping document, and the work is a multi-wave capability across the KV interface, the runner and the model, whose W3 touches every model's allocation path. Owned by row `KV-DSV4-MULTICACHE` and listed under `## Owed` in [kv-dsv4-multicache.md](specs/kv-dsv4-multicache.md) | bug | | [#1924](https://github.com/mudler/vllm.cpp/issues/1924) | `MODEL-TEXT-deepseek-v4-deepseek-v4-for-causal-lm` | **`vllm-server` refused every DeepSeek-V4-Flash SAFETENSORS checkpoint before reading a weight byte: `tokenizer: expected exactly one Split pre-tokenizer, found 3`** (`DetectPattern`, `src/vllm/tokenizer/tokenizer.cpp`). The checkpoint's `pre_tokenizer` is a FOUR-stage HF `Sequence` — `Split(\p{N}{1,3})`, `Split([一-龥぀-ゟ゠-ヿ]+)`, a `Split` over a six-alternative word regex, then `ByteLevel(use_regex=false)` — and neither reader in the tree could take it: the generic walk accepts exactly one `Split`, and `IsDeepSeekPreTokenizer` demands the SEVEN-element shape. **The two are different families and the tree's naming said otherwise**: what `SplitPattern::kDeepSeek` implements is llama.cpp's `LLAMA_VOCAB_PRE_TYPE_DEEPSEEK_LLM` (GGUF pre `deepseek-llm`, DeepSeek-V2/V2-Lite), while DeepSeek-V3, R1 and V4-Flash are `LLAMA_VOCAB_PRE_TYPE_DEEPSEEK3_LLM` (pre `deepseek-v3`) — two separate cases of one switch at `src/llama-vocab.cpp:308-325` @ `b10451`, sharing no regex. `include/vllm/tokenizer/pretokenizer.h` and `docs/reference/model-loading.md` both claimed V3 for the seven-stage row; both are corrected. FIXED by `SplitPattern::kDeepSeekV3`, a three-`Split` pipeline over the existing `ApplySplitIsolated` seam, plus a verbatim recognizer beside `IsDeepSeekPreTokenizer`. **Stage ORDER is the semantics, not a detail**: stage 2's alternation matches NO digit (`\p{N}` is in none of its six alternatives) and its rule-2 prefix class `[^\r\n\p{L}\p{P}\p{S}]?` MATCHES one, so without stage 0 first "abc123def" tokenizes as "abc" + "123def". Gated against the REAL 6.4 MB `tokenizer.json` copied byte-for-byte from the checkpoint (sha256 `8f9f37ca…33cf`), never a fixture built from a reading of the shape, with ids from HF `tokenizers` 0.22.2 — the same library vLLM's `AutoTokenizer` runs. Live differential through `examples/tokenize` on the checkpoint file: 0 id diffs over the 112-entry corpus and over 6000 random strings. The residual, recorded rather than closed: our generated tables are Unicode 15.0.0 and onig's are newer, so a codepoint UNASSIGNED at 15.0 can classify differently — pre-existing, repo-wide, and it moves no assigned codepoint (0 mismatches over 40000 random assigned-codepoint strings against 1271/30000 when unassigned ones are admitted) | bug | | [#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 | +| [#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 | | [#1957](https://github.com/mudler/vllm.cpp/issues/1957) | `LTX25-DEVICE-RESIDENCY` | **`test_ltx2_video`'s `artifacts.frames` coverage floor decides by box load, and the span-slack bound deliberately does not hold it.** Found while gating [#1946](https://github.com/mudler/vllm.cpp/issues/1946) against the three-way merge of [#1929](https://github.com/mudler/vllm.cpp/issues/1929) and [#1932](https://github.com/mudler/vllm.cpp/issues/1932); PRE-EXISTING and not that row's defect, because `git diff --name-only` across all three returns no LTX-2.5 file. `tests/vllm/multimodal/test_ltx2_video.cpp:4267` asserts `covered >= c.min_coverage * leaf_seconds` and read `CHECK( 0.000336164 >= 0.000474467 )` — a margin of **0.000138 s**, with the one part `artifacts.frames.ppm` covering 35.4254% against the 0.50 floor configured at `:4666`. NOT [#1494](https://github.com/mudler/vllm.cpp/issues/1494), which is the same assertion and the same argument but measured and CLOSED on the `denoise` leaf; closing one does not close the other, which is what #1494 itself recorded about [#1439](https://github.com/mudler/vllm.cpp/issues/1439). [#1559](https://github.com/mudler/vllm.cpp/issues/1559)'s 30 ms per-record bound does not cover this leaf and the test says so at `:4255-4259` — "(1c) does NOT hold this leaf's short records ... so for `artifacts.frames` on the nine-frame render this floor is the only thing under the head and the tail" — so the only guard is a RATIO on a sub-millisecond leaf, whose residue is two instrument boundaries plus scheduler latency while the leaf it divides into does not grow with load. MEASURED on one binary with no source change at loadavg 18-35: `-j 4` run 1 RED, `-j 4` run 2 green, `-j 2` green at `100% tests passed, 0 tests failed out of 626`, binary alone green at 106 cases / 4792 assertions. The two `-j 4` runs failed on DISJOINT tests, the other being `test_engine_core_proc`, itself recorded in `.agents/environment.md` as a parallel-ctest timing flake. Cost is attribution rather than the red, which is [#433](https://github.com/mudler/vllm.cpp/issues/433)'s class. NOT FIXED IN FLOW: bounding the residue in SECONDS beside the ratio, or gating on a quiescence precondition and reporting "not measured", changes a gate's semantics and needs its own row, spec and red-first evidence per `AGENTS.md` `## Changing the rules or a checker` — the conclusion #1439, #1494 and [#1576](https://github.com/mudler/vllm.cpp/issues/1576) each reached independently | bug | | [#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 | +| [#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 | +| [#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 | -| [#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 | +| [#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 | | [#1991](https://github.com/mudler/vllm.cpp/issues/1991) | `MODEL-MM-QWEN4-EXP` | **W4, Qwen Sparse Attention: the indexer, its side cache, and a GATHER consumer rather than a mask.** New TU `src/vllm/model_executor/models/qwen4_exp_qsa.{h,cpp}` (`indexer_n_heads` 4, `indexer_kv_heads` 1, `indexer_head_dim` 128, `indexer_budget` 2048 tokens, `indexer_compress_ratio` 4, so `block_topk` 512 and an index buffer 2051 wide). Ported onto DeepSeek-V4's C4 indexer lane per the spec's reversal, with the two things it must NOT inherit reconciled in code: DSv4's `weights_proj` + `head_scale` fold (QSA's weight is the constant `1/sqrt(128)`, applied after the sum over heads) and DSv4's GPT-J RoPE over a TRAILING span (QSA is NeoX `rotate_half` over the LEADING 64 dims with the NoPE dims trailing — the halves are swapped end for end AND the pairing convention differs). `SparseAttnCompressNormRopeStoreC4Kernel` supplies scaffolding only: its `tl.softmax(score, dim=0)` pool over an overlapping window is replaced by an unweighted mean over a non-overlapping window of 4. **The consumer is a GATHER, and that is the wave's point.** llama.cpp [#27739](https://github.com/ggml-org/llama.cpp/pull/27739) records that a sparse mask over a dense cache costs the same as dense attention under CUDA flash attention, because `flash_attn_mask_to_KV_max` only scans back to the first tile that is not all `-inf`; a mask-only QSA is CORRECT and forfeits the lever silently. Proved rather than asserted: mutating `QsaGatherAttention` to delegate to the mask reference leaves **11 of 12 cases green** and fails only `qsa-consumer: the GATHER touches only the selected rows`, so the `keys_visited` counter is the only instrument in the suite that can tell the two apart. Goldens come from RUNNING the unmodified `Qwen4ExpTextQSAIndexer.forward` at the transformers 5.16.0 lane pin (`tests/vllm/models/fixtures/gen_qwen4_exp_qsa_goldens.py`), which resolves the spec's UNMEASURED version string: `transformers.__version__` reads **5.16.0** on a live import, and the oracle runs on CPU with random weights, no checkpoint and no GPU. Both free oracles are used: sub-budget bit-identity with dense attention at kv_len = `budget + compress_ratio - 1`, and index-set equality including the ragged tail above the budget. Side cache is `MLAAttentionSpec(num_kv_heads=1, head_size=128, tokens_per_state=4)` = **64 B/token/layer at bf16**, asserted rather than defaulted, with `StatesForTokens` flooring because a ragged block writes no state. LANDS UNREACHED: `Qwen4ExpTextModel` does not exist (W2 #1987, W3 #1988, W6a #1989 are sibling waves; the registry and runner wiring is W5), row `MODEL-MM-QWEN4-EXP` owns it, [#1978](https://github.com/mudler/vllm.cpp/issues/1978) tracks it, and the spec lists it under `## Owed`. No CUDA arm and no speed claim — the speed axis opens at G4, after W6a | feature | | [#1987](https://github.com/mudler/vllm.cpp/issues/1987) | `MODEL-MM-QWEN4-EXP` | **W2 of [#1978](https://github.com/mudler/vllm.cpp/issues/1978): the hashed n-gram embedding and the PLE dilated depthwise conv, the ONLY two components of `Qwen4ExpForConditionalGeneration` with no vLLM op at all.** The negative is confirmed rather than unfound — at vLLM `origin/main` = `6a5e8f5979`, `git grep -in dilat` returns ZERO lines in `layers/mamba/`, zero in `csrc/` and zero in `tests/`, and `layers/conv.py` defines only `Conv2dLayer`/`Conv3dLayer` — so transformers **v5.16.0**, this row's accepted lane pin, is the sole oracle. Three silent divergence sites, all verified: `_splitmix64` must be `uint64_t` throughout (its `>> 30/27/31` are LOGICAL shifts on a non-negative Python int; on `int64_t` they go arithmetic and the multiplicand is wrong, with the top bit set about half the time), `_splitmix64(value) % half_bound` must be an UNSIGNED modulo (the dividend routinely exceeds 2^63), and shard reassembly is NUMERIC not lexicographic (`shard_0, shard_1, shard_10, ...` silently permutes a 95 GiB table). The forward is int64-exact and needs no bignum because `multiplier_max * vocab_size <= 2^63 - 1` — but **that bound holds only while every token id is below `vocab_size`**, so an out-of-range id overflows and diverges in silence, and this wave refuses one by name. Goldens confirmed a FOURTH way beyond #1987's three: `vocab_size = 248320`, read from the released `config.json`, is the UNIQUE preimage below 2e6 of the published `layer_multipliers = [23703573157769, 20109073645365, 8052911324071]` at the absent-`seed` default 1234. PLE specifics: the conv state is **9 columns** — `(kernel-1) * dilation`, NOT `kernel-1` — and holds the **normed** conv input while the skip term is the **un-normed** copy; the signed-sqrt gate clamps BEFORE the sqrt so the output magnitude floor is `sqrt(1e-6) = 1e-3` and tiny scores are AMPLIFIED, with exactly zero mapping to zero, making the function genuinely discontinuous at an origin that a fully masked row reaches; and the n-gram history must be padded with EOS because upstream's `update_conv_state` pads with 0, a valid token id. **Lands UNREACHED** per AGENTS.md "Nothing lands dead": no production entry point calls it until W5 assembles the model, the wiring is owned by `MODEL-MM-QWEN4-EXP` and tracked by #1978, and the spec lists it under `## Owed`. Host reference only; the batched device arm is owed, and the per-sequence signatures exist so it drops in — the gather is 16 uncoalesced random rows per token PER SEQUENCE, so at batch B it is 16*B independent gathers with no coalescing. No token claim, no speed claim: `## Gates` admits none from this row until an arm runs, and nothing published fits any fleet device | feature | -| [#1963](https://github.com/mudler/vllm.cpp/issues/1963) | `ROAD-V1-MEM` | At ctx=32768 `--max-num-seqs 32` our engine consumes ~108 GB during load and never serves; vLLM and SGLang both serve there. Root cause found and fixed by `FIX-KV-GROUP-LAYER-COUNT` ([spec](../specs/kv-group-layer-count.md)): thirty-three of thirty-four registries publish ONE placeholder name per KV group, `KVBytesPerBlock` reads `layer_names.size()` as the layer count, and `ResolveNumBlocks` arm 2 therefore divides an absolute `--kv-cache-memory` budget by ONE layer's page while the runner allocates one buffer per layer — measured 8.5 GiB allocated for a 1 GiB budget on the 27B | bug | -| [#1966](https://github.com/mudler/vllm.cpp/issues/1966) | `ROAD-V1-MEM` | The #371 recurrent-state OOM guard is 48x under and passes a config that allocates 43.4 GiB, because `recurrent_state_bytes` counts placeholder layer names. Same root defect as [#1963](https://github.com/mudler/vllm.cpp/issues/1963), different code path; both fixed by `FIX-KV-GROUP-LAYER-COUNT` ([spec](../specs/kv-group-layer-count.md)) | bug | -| [#1982](https://github.com/mudler/vllm.cpp/issues/1982) | `SERVE-STREAM-USAGE` | **`ChatSseStream::next` writes the `/v1/chat/completions` role frame before it reads anything from the engine, so `vllm bench serve --backend openai-chat` stamps TTFT on an empty frame and our TTFT through that harness is an HTTP round trip, not a time to first token.** Upstream builds the role chunk under `if first_iteration:` inside `async for res in result_generator:` (`vllm/entrypoints/openai/chat_completion/serving.py:477,487`) and says why at `:484-486`: an exception in the generator "needs to be sent as the FIRST response". `vllm/benchmarks/lib/endpoint_request_func.py:404-408` guards on the presence of `choices`, not on non-empty `delta.content`, and our role frame carries `delta.content = ""` with no `usage`. vLLM and SGLang order the frame after the first result, so their rows on the same harness are honest and only ours is not; this blocks the #1574 three-engine TTFT row. `.agents/specs/stream-options.md` scoped the buffering to continuous usage on purpose and both its passages are corrected here. Fixed by removing the `usage_.include_continuous_usage` guard around the first-result buffering loop, so the default path buffers too. Spec: [`specs/chat-role-frame-ordering.md`](../specs/chat-role-frame-ordering.md) | bug | -| [#1992](https://github.com/mudler/vllm.cpp/issues/1992) | — | **Neither `ChatSseStream::next` nor `CompletionSseStream::next` converts an engine exception into a `data: {"error": …}` frame, so a streaming request that fails is a truncated 200 and the cause reaches only `stderr`.** Upstream yields the error frame and then `data: [DONE]` from the generator's `except GenerationError` / `except Exception` arms (`vllm/entrypoints/openai/chat_completion/serving.py:827-833` at the pin `555967922`), and that frame is what makes the first-iteration ordering at `:484-486` mean anything: the role chunk is built inside the loop so an exception can be the FIRST response, which needs a response to exist. Ours propagates out of `next()` into the cpp-httplib chunked content provider (`src/vllm/entrypoints/openai/api_server.cpp::ApiServer::register_routes`), which logs `sse: stream aborted mid-flight:` and aborts, so a client cannot tell a failed request from a short one. Found while fixing [#1982](https://github.com/mudler/vllm.cpp/issues/1982) and NOT fixed in that flow: upstream's `try` wraps the whole generator, so the frame is owed for mid-stream failures on both endpoints, and that is a different blast radius needing its own red-first cases for the payload shape, the trailing `[DONE]` and the separate `GenerationError` converter. Owed by [`specs/chat-role-frame-ordering.md`](../specs/chat-role-frame-ordering.md) `## Owed` | bug | -| [#1983](https://github.com/mudler/vllm.cpp/issues/1983) | `KV-GDN-STATE-BUDGET` | **The GDN recurrent-state pool is preallocated per CONFIGURED sequence, on an axis no flag bounds.** `GPUModelRunner::initialize_kv_cache` sizes `gdn_state_slots_ = max_num_reqs * (num_spec + 1)` and allocates one conv and one SSM buffer per GDN layer from it, each `Memset` to zero at construction, so every byte is resident before the first request. Re-derived for `Qwen3.8-27B` (48 linear-attention layers, `Hk/Hv/Dk/Dv/conv = 16/48/128/128/4`, `mamba_ssm_dtype = float32`) at `num_speculative_tokens = 8`: one slot costs 3,371,008 B per layer, 154.31 MiB across 48 layers, so one sequence costs 1.356 GiB and `--max-num-seqs 32` costs **43.40 GiB** that `--kv-cache-memory`, `--num-blocks` and `--gpu-memory-utilization` all fail to bound. The per-sequence cost is NOT the divergence — upstream charges the same `1 + num_speculative_blocks` state blocks (`vllm/v1/kv_cache_interface.py::MambaSpec.max_memory_usage_bytes`) and our `f32` SSM mirrors the checkpoint's own `mamba_ssm_dtype` — the AXIS is: `max_num_seqs` sizes no allocation anywhere in vLLM. Upstream raises the attention block size until one attention page holds one mamba page (`vllm/platforms/interface.py::Platform.check_and_update_config`), pads the mamba page to match, and then draws BOTH from one budgeted pool whose tensors are `shared_by` one layer from each group (`kv_cache_utils.py::_get_kv_cache_config_uniform_page_size`), so its recurrent allocation is a function of available memory and never of the concurrency cap. Fixed by mirroring that arithmetic in `ComputeHybridKvBudget` — `unified_block_tokens = align * cdiv(mamba_page, align * attn_bytes_per_token)`, `max_state_seqs = (num_blocks * block_size / unified_block_tokens) / (1 + num_spec)` — and resolving ONE `max_num_seqs` from it for the runner, the scheduler and the #371 guard alike. The bound reads no layer count (upstream's per-layer page equality cancels it), so it does not depend on the placeholder-layer-name repair owned by [#1963](https://github.com/mudler/vllm.cpp/issues/1963) and [#1966](https://github.com/mudler/vllm.cpp/issues/1966), and it lands in its own translation unit so the three rows share no edit surface. Spec: [`specs/gdn-state-kv-budget.md`](../specs/gdn-state-kv-budget.md) | bug | -| [#2021](https://github.com/mudler/vllm.cpp/issues/2021) | `BACKEND-TENSTORRENT-QWEN35` | `qwen3_5.cpp`'s Tenstorrent debug-dump path calls `std::fwrite(vec.data(), ...)` on a `std::vector` that may be empty, and GCC 15.2.0 rejects the call statically under `-Werror=nonnull` because `fwrite`'s first parameter carries the `nonnull` attribute — a size-0 write on a null pointer is well defined by the C standard, but the attribute makes the call itself UB by its letter, so every fresh CPU build on that toolchain fails. Fixed in flow by guarding the write on `!vec.empty()`, mirroring the raw-bytes write three lines above | bug | +| [#1963](https://github.com/mudler/vllm.cpp/issues/1963) | `ROAD-V1-MEM` | At ctx=32768 `--max-num-seqs 32` our engine consumes ~108 GB during load and never serves; vLLM and SGLang both serve there. Root cause found and fixed by `FIX-KV-GROUP-LAYER-COUNT` ([spec](specs/kv-group-layer-count.md)): thirty-three of thirty-four registries publish ONE placeholder name per KV group, `KVBytesPerBlock` reads `layer_names.size()` as the layer count, and `ResolveNumBlocks` arm 2 therefore divides an absolute `--kv-cache-memory` budget by ONE layer's page while the runner allocates one buffer per layer — measured 8.5 GiB allocated for a 1 GiB budget on the 27B | bug | +| [#1966](https://github.com/mudler/vllm.cpp/issues/1966) | `ROAD-V1-MEM` | The #371 recurrent-state OOM guard is 48x under and passes a config that allocates 43.4 GiB, because `recurrent_state_bytes` counts placeholder layer names. Same root defect as [#1963](https://github.com/mudler/vllm.cpp/issues/1963), different code path; both fixed by `FIX-KV-GROUP-LAYER-COUNT` ([spec](specs/kv-group-layer-count.md)) | bug | +| [#1982](https://github.com/mudler/vllm.cpp/issues/1982) | `SERVE-STREAM-USAGE` | **`ChatSseStream::next` writes the `/v1/chat/completions` role frame before it reads anything from the engine, so `vllm bench serve --backend openai-chat` stamps TTFT on an empty frame and our TTFT through that harness is an HTTP round trip, not a time to first token.** Upstream builds the role chunk under `if first_iteration:` inside `async for res in result_generator:` (`vllm/entrypoints/openai/chat_completion/serving.py:477,487`) and says why at `:484-486`: an exception in the generator "needs to be sent as the FIRST response". `vllm/benchmarks/lib/endpoint_request_func.py:404-408` guards on the presence of `choices`, not on non-empty `delta.content`, and our role frame carries `delta.content = ""` with no `usage`. vLLM and SGLang order the frame after the first result, so their rows on the same harness are honest and only ours is not; this blocks the #1574 three-engine TTFT row. `.agents/specs/stream-options.md` scoped the buffering to continuous usage on purpose and both its passages are corrected here. Fixed by removing the `usage_.include_continuous_usage` guard around the first-result buffering loop, so the default path buffers too. Spec: [`specs/chat-role-frame-ordering.md`](specs/chat-role-frame-ordering.md) | bug | +| [#1992](https://github.com/mudler/vllm.cpp/issues/1992) | — | **Neither `ChatSseStream::next` nor `CompletionSseStream::next` converts an engine exception into a `data: {"error": …}` frame, so a streaming request that fails is a truncated 200 and the cause reaches only `stderr`.** Upstream yields the error frame and then `data: [DONE]` from the generator's `except GenerationError` / `except Exception` arms (`vllm/entrypoints/openai/chat_completion/serving.py:827-833` at the pin `555967922`), and that frame is what makes the first-iteration ordering at `:484-486` mean anything: the role chunk is built inside the loop so an exception can be the FIRST response, which needs a response to exist. Ours propagates out of `next()` into the cpp-httplib chunked content provider (`src/vllm/entrypoints/openai/api_server.cpp::ApiServer::register_routes`), which logs `sse: stream aborted mid-flight:` and aborts, so a client cannot tell a failed request from a short one. Found while fixing [#1982](https://github.com/mudler/vllm.cpp/issues/1982) and NOT fixed in that flow: upstream's `try` wraps the whole generator, so the frame is owed for mid-stream failures on both endpoints, and that is a different blast radius needing its own red-first cases for the payload shape, the trailing `[DONE]` and the separate `GenerationError` converter. Owed by [`specs/chat-role-frame-ordering.md`](specs/chat-role-frame-ordering.md) `## Owed` | bug | +| [#1983](https://github.com/mudler/vllm.cpp/issues/1983) | `KV-GDN-STATE-BUDGET` | **The GDN recurrent-state pool is preallocated per CONFIGURED sequence, on an axis no flag bounds.** `GPUModelRunner::initialize_kv_cache` sizes `gdn_state_slots_ = max_num_reqs * (num_spec + 1)` and allocates one conv and one SSM buffer per GDN layer from it, each `Memset` to zero at construction, so every byte is resident before the first request. Re-derived for `Qwen3.8-27B` (48 linear-attention layers, `Hk/Hv/Dk/Dv/conv = 16/48/128/128/4`, `mamba_ssm_dtype = float32`) at `num_speculative_tokens = 8`: one slot costs 3,371,008 B per layer, 154.31 MiB across 48 layers, so one sequence costs 1.356 GiB and `--max-num-seqs 32` costs **43.40 GiB** that `--kv-cache-memory`, `--num-blocks` and `--gpu-memory-utilization` all fail to bound. The per-sequence cost is NOT the divergence — upstream charges the same `1 + num_speculative_blocks` state blocks (`vllm/v1/kv_cache_interface.py::MambaSpec.max_memory_usage_bytes`) and our `f32` SSM mirrors the checkpoint's own `mamba_ssm_dtype` — the AXIS is: `max_num_seqs` sizes no allocation anywhere in vLLM. Upstream raises the attention block size until one attention page holds one mamba page (`vllm/platforms/interface.py::Platform.check_and_update_config`), pads the mamba page to match, and then draws BOTH from one budgeted pool whose tensors are `shared_by` one layer from each group (`kv_cache_utils.py::_get_kv_cache_config_uniform_page_size`), so its recurrent allocation is a function of available memory and never of the concurrency cap. Fixed by mirroring that arithmetic in `ComputeHybridKvBudget` — `unified_block_tokens = align * cdiv(mamba_page, align * attn_bytes_per_token)`, `max_state_seqs = (num_blocks * block_size / unified_block_tokens) / (1 + num_spec)` — and resolving ONE `max_num_seqs` from it for the runner, the scheduler and the #371 guard alike. The bound reads no layer count (upstream's per-layer page equality cancels it), so it does not depend on the placeholder-layer-name repair owned by [#1963](https://github.com/mudler/vllm.cpp/issues/1963) and [#1966](https://github.com/mudler/vllm.cpp/issues/1966), and it lands in its own translation unit so the three rows share no edit surface. Spec: [`specs/gdn-state-kv-budget.md`](specs/gdn-state-kv-budget.md) | bug | | [#2018](https://github.com/mudler/vllm.cpp/issues/2018) | `ENG-HYBRID-PLACEMENT` | W1: the `placement` sub-object of `--offload-config`'s `vllm_cpp` key, mapping llama.cpp's four placement surfaces at pin `b10451` — `overrides` (`-ot`) as the general form, `cpu_moe` (`-cmoe`) and `n_cpu_moe` (`-ncmoe N`) as sugar that desugars into it, and `fit` (`--fit`) refused beside a manual placement. Config only: it parses, validates, refuses, merges, resolves precedence and desugars, and no weight or compute moves | feature | | [#2023](https://github.com/mudler/vllm.cpp/issues/2023) | `ENG-HYBRID-PLACEMENT` | W2: `DevicePlacement`, the resolved answer to which device runs a tensor — first-match-wins over the operator's ordered override list, `regex_search` rather than a full match, regexes compiled once at build, and an `IsTrivial()` inertness predicate that is true both for no overrides and for overrides naming the engine's own device. Resolves and REPORTS; W3 owns the routing that reads it | feature | | [#2012](https://github.com/mudler/vllm.cpp/issues/2012) | `ENG-HYBRID-PLACEMENT` | `docs/FEATURES.md` compares eleven memory axes against llama.cpp and carries no row for hybrid CPU/GPU expert placement, so `-cmoe`/`-ncmoe` — a capability this engine does not have, owned by the `READY` row `ENG-HYBRID-PLACEMENT` and requested in #149 — is invisible in the comparison. The nearest row, routed-expert streaming from disk, records llama.cpp as `mmap only`, which is correct for that row and is why the gap hides: streaming moves weights toward the compute, placement moves compute toward the weights | record | | [#2015](https://github.com/mudler/vllm.cpp/issues/2015) | `ENG-HYBRID-PLACEMENT` | Map llama.cpp's four placement surfaces (`-ot` general, `-cmoe` and `-ncmoe N` as sugar over it, `--fit` auto) onto the `vllm_cpp` extension key of `--offload-config` instead of a bespoke flag, per developer direction 2026-08-26, and re-anchor the spec at the recorded pin `b10451` — `237ad9b96` was superseded on 2026-08-22 and none of its cited line numbers resolve | record | | [#2005](https://github.com/mudler/vllm.cpp/issues/2005) | `BACKEND-TENSTORRENT` | **No Tenstorrent sibling of `gpu_clock_state.py` existed, so every TT speed figure on record was clock-unattributed — including #2003's inversion, which could have been a clock excursion.** Landed `tools/bench/tt_clock_state.py` (+`tt_refold_busy.py`): AICLK via 430 ms `tt-smi -s` snapshots at 1 Hz, NVIDIA thresholds copied constant-for-constant (spread 5%, cross-arm median AND mean 1%, 30 busy samples, majority-busy), stop-only summaries (#1657 shape), NOT-APPLICABLE fields stated (no persistence/applications knobs; no live throttle bitmap → carried caveat), claimed-max as provenance-carrying argument, busy proxied by the leg pid holding `/dev/tenstorrent` fds checked live per interval, machine identity compared unconditionally even under a boot waiver. First wired use re-measured #2003 same-day: the P150 AICLK governor proved TWO-STATE (800 idle / pegged cap under load); raw windows refuse within-run spread 40.74% by construction of that platform, and the live-recorded busy flags enable a criterion-independent busy-slice refold in which **all six windows carry exactly one distinct value {1350}, spread 0.00%, judge PASS** — the inverted ratio is clock-attributed parity. Owed: verified claimed-max pin (1350 is class folklore today), in-process pyluwen sampler for sub-second cadence, and a policy decision on whether two-state governors deserve a spread rule scoped to busy slices inside the tool itself rather than the offline refold script | perf | -| [#1954](https://github.com/mudler/vllm.cpp/issues/1954) | — | **`test_backend_cross_device` is red on gfx1200 before `ROCM-KQUANT-NWARPS-DECODE` and after it**, at `CHECK(got == ref_b)` in the case "MoeSiluMul matches the CPU oracle within NMSE <= 5e-4", bf16 arm, 1 of 26 cases and 1 of 80253 assertions, on an RX 9060 XT under ROCm 7.2.3. Elements differ in the last digit. It is the gfx1200 counterpart of the CUDA-only [#1802](https://github.com/mudler/vllm.cpp/issues/1802) and [#907](https://github.com/mudler/vllm.cpp/issues/907), which record the same test name and the same assertion on sm_110 and sm_121a. **Proven pre-existing rather than assumed**: reverting that row's two source files to the parent `5888abf11` and rebuilding reproduces the identical failure at 24 of 25 cases and 1 of 80195 assertions, and the control was run twice, by two independent reviews. NOT FIXED IN FLOW: nothing in the `KQuantGemmK` decode path that row touches reaches `MoeSiluMul`, and a last-digit numerics defect in a fused activation needs its own row, spec and red-before evidence rather than a repair folded into an `nwarps` sweep. `.agents/environment.md` now carries a gfx1200 known-red table, keyed on the test name the way the sm_110 table is keyed, so the next reader of that gate reads one expected red rather than a regression. The line number is a pointer to re-derive and not part of that key: the issue recorded the assertion at `:2063` and the row's `CAPTURE` stringification repair moved it to `:2067`. Listed under `## Owed` in [`rocm-kquant-nwarps-decode.md`](../specs/rocm-kquant-nwarps-decode.md) | bug | -| [#2013](https://github.com/mudler/vllm.cpp/issues/2013) | `BACKEND-ROCM` | **RETRACTS the #907 half of the [#1954](https://github.com/mudler/vllm.cpp/issues/1954) row directly above.** That row calls the gfx1200 red "the gfx1200 counterpart of the CUDA-only [#1802](https://github.com/mudler/vllm.cpp/issues/1802) and [#907](https://github.com/mudler/vllm.cpp/issues/907), which record the same test name and the same assertion on sm_110 and sm_121a". **#1802 does. #907 does NOT.** #1954 is `test_backend_cross_device`'s "MoeSiluMul matches the CPU oracle within NMSE <= 5e-4" case, at `:2063` `CHECK(got == ref_b)`, bf16 arm, on an RX 9060 XT under ROCm 7.2.3 on gfx1200. [#907](https://github.com/mudler/vllm.cpp/issues/907) is `test_cuda_ops` at 439 of 440 assertions on GB10, sm_121a: a different binary, a different case, and a different assertion. Both are last-digit bf16 numerics in a silu-and-mul kernel, so the two share a DEFECT FAMILY and nothing narrower. #907 is not a second recorded instance of #1954's test or of its assertion. **The case attribution is thinner than the corrected text reads.** #907's own body names the binary and the count only, and never names a case. The case "CUDA silu_and_mul matches CPU" (`tests/vt/test_cuda_ops.cpp:330`) is read from the sm_110 table row at `.agents/environment.md:1332` and from the [#1802](https://github.com/mudler/vllm.cpp/issues/1802) row, which pair GB10's 439 of 440 against sm_110's 438 of 440 in that one case. That is an inference across two records, not a GB10 log naming the failing case, and the line number is a pointer to re-derive rather than a key. **Why a whole row and not an edit**: this index is append-only under `merge=union`, and `scripts/check-agent-record.py` refuses a second row bearing a number the index already carries, so the #1954 row can be corrected neither in place nor by re-citing #1954. The correction needed a fresh number, and this issue is that number. The two other in-repo copies were repaired at `e76b18463`: the gfx1200 known-red row in [`environment.md`](../environment.md) and the `## Owed` entry in [`rocm-kquant-nwarps-decode.md`](../specs/rocm-kquant-nwarps-decode.md). #1954's own GitHub title and body were corrected before that, and the corrected body is the source of truth for this distinction. Same shape as [#1339](https://github.com/mudler/vllm.cpp/issues/1339) superseding [#1280](https://github.com/mudler/vllm.cpp/issues/1280) and [#1796](https://github.com/mudler/vllm.cpp/issues/1796) retracting [#1456](https://github.com/mudler/vllm.cpp/issues/1456). The gfx1200 red itself does not change: #1954 stays open, stays listed under `## Owed` in [`rocm-kquant-nwarps-decode.md`](../specs/rocm-kquant-nwarps-decode.md), and this row corrects only what that red is a counterpart OF | record | | [#2032](https://github.com/mudler/vllm.cpp/issues/2032) | `ENG-HYBRID-PLACEMENT` | W3a: `MoePlacementPlan`, resolving a name-keyed `DevicePlacement` to the per-layer decision the forward can actually read, against llama.cpp's own GGUF tensor spelling — and REFUSING a partial placement, where a layer's gate, up and down do not share a device, because the MoE block runs one grouped GEMM over the three | feature | -| [#2029](https://github.com/mudler/vllm.cpp/issues/2029) | `ENG-CUDAGRAPH-BREAK` | With DFlash2 speculation OFF the engine dies at c=8 in CUDA graph capture: `cudaMalloc: operation not permitted when stream is capturing`. Located statically: `Pool(b).PreGrowForCapture(b, s.demand)` — the #1380 capture pre-grow — sits INSIDE `if (dbuf)` in both Qwen3.5 decode-graph drivers (`qwen3_5.cpp:10885/10907`, `:11439/11461`), and `dbuf = impl_->dbuf \|\| spec_step` is false on the DEFAULT server, where `VT_ASYNC_EXECUTOR` is unset and no step is speculative. So the pre-grow is exactly the "path taken only when speculation is off" the issue names, by its absence. Spec [cudagraph-pregrow-nonspec.md](../specs/cudagraph-pregrow-nonspec.md) | bug | -| [#2035](https://github.com/mudler/vllm.cpp/issues/2035) | `ENG-CUDAGRAPH-BREAK` | Seven decode-graph drivers (`qwen3.cpp`, `qwen3_moe.cpp`, `deepseek_v2.cpp`, `deepseek_v4.cpp`, `voxtral.cpp`, `laguna.cpp`, `qwen3_dflash.cpp`) open a `vt::GraphCaptureScope` with no `DevicePool::PreGrowForCapture` and no demand profile at all — only `qwen3_5.cpp` uses any of the #1380 machinery. Found while fixing #2029; owed under `## Owed` in [cudagraph-pregrow-nonspec.md](../specs/cudagraph-pregrow-nonspec.md) | bug | -| [#2036](https://github.com/mudler/vllm.cpp/issues/2036) | `ENG-CUDAGRAPH-BREAK` | `DenseAlignFor` (`qwen3_5.cpp:2825-2849`) allocates five blocks and calls `d.b.Synchronize(d.q)` at `:2846` on an `M` miss, and `EnsureCtmp` (`cuda_marlin_dense.cu:74-89`) grows with `cudaMallocAsync` at `:85`, both with no `cudaStreamIsCapturing` refusal — unlike the six sibling shape-keyed caches that have one. Latent today (the cold step visits the same key), found while fixing #2029; owed under `## Owed` in [cudagraph-pregrow-nonspec.md](../specs/cudagraph-pregrow-nonspec.md) | bug | -| [#2037](https://github.com/mudler/vllm.cpp/issues/2037) | — | `EngineDeadError` promises "See stack trace (above)" (`include/vllm/v1/engine/core_client.h:63`) and the fatal handler prints only `e.what()` (`src/vllm/v1/engine/core_client.cpp:36-38`), so no trace is ever emitted. #1380 closed only because somebody instrumented `CudaBackend::Alloc` by hand; #2028 and #2029 both record the gap. Owed under `## Owed` in [cudagraph-pregrow-nonspec.md](../specs/cudagraph-pregrow-nonspec.md) | bug | +| [#2029](https://github.com/mudler/vllm.cpp/issues/2029) | `ENG-CUDAGRAPH-BREAK` | With DFlash2 speculation OFF the engine dies at c=8 in CUDA graph capture: `cudaMalloc: operation not permitted when stream is capturing`. Located statically: `Pool(b).PreGrowForCapture(b, s.demand)` — the #1380 capture pre-grow — sits INSIDE `if (dbuf)` in both Qwen3.5 decode-graph drivers (`qwen3_5.cpp:10885/10907`, `:11439/11461`), and `dbuf = impl_->dbuf \|\| spec_step` is false on the DEFAULT server, where `VT_ASYNC_EXECUTOR` is unset and no step is speculative. So the pre-grow is exactly the "path taken only when speculation is off" the issue names, by its absence. Spec [cudagraph-pregrow-nonspec.md](specs/cudagraph-pregrow-nonspec.md) | bug | +| [#2035](https://github.com/mudler/vllm.cpp/issues/2035) | `ENG-CUDAGRAPH-BREAK` | Seven decode-graph drivers (`qwen3.cpp`, `qwen3_moe.cpp`, `deepseek_v2.cpp`, `deepseek_v4.cpp`, `voxtral.cpp`, `laguna.cpp`, `qwen3_dflash.cpp`) open a `vt::GraphCaptureScope` with no `DevicePool::PreGrowForCapture` and no demand profile at all — only `qwen3_5.cpp` uses any of the #1380 machinery. Found while fixing #2029; owed under `## Owed` in [cudagraph-pregrow-nonspec.md](specs/cudagraph-pregrow-nonspec.md) | bug | +| [#2036](https://github.com/mudler/vllm.cpp/issues/2036) | `ENG-CUDAGRAPH-BREAK` | `DenseAlignFor` (`qwen3_5.cpp:2825-2849`) allocates five blocks and calls `d.b.Synchronize(d.q)` at `:2846` on an `M` miss, and `EnsureCtmp` (`cuda_marlin_dense.cu:74-89`) grows with `cudaMallocAsync` at `:85`, both with no `cudaStreamIsCapturing` refusal — unlike the six sibling shape-keyed caches that have one. Latent today (the cold step visits the same key), found while fixing #2029; owed under `## Owed` in [cudagraph-pregrow-nonspec.md](specs/cudagraph-pregrow-nonspec.md) | bug | +| [#2037](https://github.com/mudler/vllm.cpp/issues/2037) | — | `EngineDeadError` promises "See stack trace (above)" (`include/vllm/v1/engine/core_client.h:63`) and the fatal handler prints only `e.what()` (`src/vllm/v1/engine/core_client.cpp:36-38`), so no trace is ever emitted. #1380 closed only because somebody instrumented `CudaBackend::Alloc` by hand; #2028 and #2029 both record the gap. Owed under `## Owed` in [cudagraph-pregrow-nonspec.md](specs/cudagraph-pregrow-nonspec.md) | bug | | [#2026](https://github.com/mudler/vllm.cpp/issues/2026) | `ENG-HYBRID-PLACEMENT` | W3b: route a placed layer's routed-expert compute to the placement device, with the activation round trip at the boundary — and refuse the fp4-resident arm, whose device residents are built eagerly at load so placing it would upload every expert and then compute across the bus, a defect a token gate cannot see | feature | | [#2040](https://github.com/mudler/vllm.cpp/issues/2040) | `ENG-HYBRID-PLACEMENT` | W3c: replace W3b's hand-written per-model round trip with ONE shared seam every architecture routes through, exploiting that every MoE block already has the same `(Dev, weights, params, [T,H] dh, T)` shape — and correct `docs/ENVIRONMENT.md` and `docs/FEATURES.md`, which W3b made false by leaving them saying the knobs place nothing | bug | | [#2049](https://github.com/mudler/vllm.cpp/issues/2049) | `MODEL-MM-QWEN4-EXP` | **The row spec's settled-decisions list stated the REFUTED QSA mapping**, telling a fresh W4/W5 implementer to build QSA on MiniMax-M3 and calling the DeepSeek-V4 lane "the wrong port" — the exact reverse of the correction recorded in the same file's Port map and Design section, in the matrix row, and in [#1978](https://github.com/mudler/vllm.cpp/issues/1978). Pre-existing on `main`; found while reviewing the W6a merge ([#2019](https://github.com/mudler/vllm.cpp/pull/2019)) and fixed in that same flow per AGENTS.md "Every change starts from an issue". Load-bearing rather than cosmetic: item 2 sits in the section written so an implementer does NOT re-derive it, and a top-down reader hits the stale instruction before the corrected Design section. | doc | | [#2046](https://github.com/mudler/vllm.cpp/issues/2046) | `ENG-HYBRID-PLACEMENT` | `qwen3_5.cpp` kept private `Dev`/`DBuf`/`MakeTensor`/`Reshape` copies instead of the shared `dense_device_glue.h` set — the off-framework divergence its own `ResidentWeight` comment records, where a repair reached 25 model files and not this one. The private types also had INTERNAL LINKAGE, which is what forced the MoE placement seam to carry a glue-templated second spelling; migrating collapses it back to one | bug | | [#2050](https://github.com/mudler/vllm.cpp/issues/2050) | `ENG-HYBRID-PLACEMENT` | Laguna's FFN is host-orchestrated token-at-a-time — per-token host rows, the router on the host through `MatmulNK`, and a host scalar combine loop — so a device-shaped MoE entry wrapping those loops would put it in the placement seam's wired list while moving nothing and adding a round trip: supported to read, a regression to measure. The real repair is a device-resident batched FFN, which is a model rework with a performance gate | gap | -| [#1984](https://github.com/mudler/vllm.cpp/issues/1984) | `SAMPLE-CORE` | `RandomSampleKernel` is launched `<<>>` and scans a 248,320-wide vocab on one thread per row, computing two `SplitMix64` rounds and an f64 `log` per element. Eleven lines above it the same file records that a single-block single-thread scan of a ~151k vocab cost ~7.5 ms/token, which is why greedy argmax was rewritten into `ArgmaxPartialKernel`/`ArgmaxFinalKernel`; the Gumbel draw never got that treatment. Upstream is whole-tensor (`vllm/v1/sample/ops/topk_topp_sampler.py::sample_with_exponential_noise`), so this is a mirror obligation. Reached by every non-greedy row through `ModelRunner::execute_model` -> `Sampler::forward` -> `vt::RandomSample`. Spec: [sample-gen-config-and-parallel-gumbel.md](../specs/sample-gen-config-and-parallel-gumbel.md) | perf | -| [#1985](https://github.com/mudler/vllm.cpp/issues/1985) | `SAMPLE-CORE` | `generation_config.json` is read for `eos_token_id` only (`hf_config.cpp::ReadGenerationConfigEosIds`), so `Qwen/Qwen3.8-27B`'s shipped `top_k: 20` / `top_p: 0.95` never reach `SamplingParams` and `to_sampling_params` resolves omitted knobs straight to the neutral OpenAI defaults, which disable both filters. vLLM applies them through `ModelConfig.get_diff_sampling_param` -> `OpenAIServing*.default_sampling_params` -> `to_sampling_params`. Since `vllm bench serve` stopped sending `--temperature`, both engines sample at temperature 1.0 and vLLM draws from 20 candidates while we draw from 248,320: different sampling on two sides of a parity benchmark. Spec: [sample-gen-config-and-parallel-gumbel.md](../specs/sample-gen-config-and-parallel-gumbel.md) | bug | +| [#1984](https://github.com/mudler/vllm.cpp/issues/1984) | `SAMPLE-CORE` | `RandomSampleKernel` is launched `<<>>` and scans a 248,320-wide vocab on one thread per row, computing two `SplitMix64` rounds and an f64 `log` per element. Eleven lines above it the same file records that a single-block single-thread scan of a ~151k vocab cost ~7.5 ms/token, which is why greedy argmax was rewritten into `ArgmaxPartialKernel`/`ArgmaxFinalKernel`; the Gumbel draw never got that treatment. Upstream is whole-tensor (`vllm/v1/sample/ops/topk_topp_sampler.py::sample_with_exponential_noise`), so this is a mirror obligation. Reached by every non-greedy row through `ModelRunner::execute_model` -> `Sampler::forward` -> `vt::RandomSample`. Spec: [sample-gen-config-and-parallel-gumbel.md](specs/sample-gen-config-and-parallel-gumbel.md) | perf | +| [#1985](https://github.com/mudler/vllm.cpp/issues/1985) | `SAMPLE-CORE` | `generation_config.json` is read for `eos_token_id` only (`hf_config.cpp::ReadGenerationConfigEosIds`), so `Qwen/Qwen3.8-27B`'s shipped `top_k: 20` / `top_p: 0.95` never reach `SamplingParams` and `to_sampling_params` resolves omitted knobs straight to the neutral OpenAI defaults, which disable both filters. vLLM applies them through `ModelConfig.get_diff_sampling_param` -> `OpenAIServing*.default_sampling_params` -> `to_sampling_params`. Since `vllm bench serve` stopped sending `--temperature`, both engines sample at temperature 1.0 and vLLM draws from 20 candidates while we draw from 248,320: different sampling on two sides of a parity benchmark. Spec: [sample-gen-config-and-parallel-gumbel.md](specs/sample-gen-config-and-parallel-gumbel.md) | bug | | [#2002](https://github.com/mudler/vllm.cpp/issues/2002) | `SAMPLE-CORE` | With `--speculative-config` set, `GPUModelRunner::sample_tokens` branches on `num_draft_tokens > 0` alone and returns the greedy-only `RejectionSampler`'s output, so `Sampler::forward` and `vt::RandomSample` are never called and a `temperature: 1.0` request decodes GREEDILY. `include/vllm/v1/spec_decode/rejection_sampler.h` states the contract it violates in its own deferral list ("a temperature > 0 request must NOT be routed here yet"); neither the runner nor `RejectionSampler::forward` enforces it. Found while writing #1984's acceptance measurement against a baseline recipe carrying `--speculative-config`, where the sampler under test would never have been launched and the null result would have read as "the change did nothing" | bug | -| [#2059](https://github.com/mudler/vllm.cpp/issues/2059) | — | `.github/workflows/ci.yml:1598` sets `VT_POOL_BYPASS: "1"` for BOTH `sanitize-cpu` lanes, so the `DevicePool` free list, size-class ladder, best-fit borrow (#1922) and capture pre-grow (#1380) are unexecuted under ASan AND TSan. The stated justification is ASan's `detect_leaks`; ThreadSanitizer has no leak detector and gains nothing. MEASURED: under `-DVLLM_CPP_SANITIZE=thread` with the pool ENABLED, `test_qwen3_5_decode_graph_seam` is 10/10, 156 assertions, exit 0, zero TSan warnings. Found while repairing the #2047 red; owed under `## Owed` in [cudagraph-pregrow-nonspec.md](../specs/cudagraph-pregrow-nonspec.md) | bug | -| [#2055](https://github.com/mudler/vllm.cpp/issues/2055) | `ENG-UPSTREAM-LTX2-PIN` | `tools/oracle/ltx2_oracle.py` asserts the LTX-2 oracle's revision and its resolved `ltx_core` / `ltx_pipelines` origins in the PARENT, then renders in a child started with `python -m`, which puts the CURRENT WORKING DIRECTORY on that child's `sys.path[0]`. The parent's `importlib.util.find_spec` never consults it, so a directory holding a decoy `ltx_pipelines`, made the CWD, is imported by the process that loads the weights while the process that checked identity sees nothing. MEASURED by a fresh reviewer of [#2053](https://github.com/mudler/vllm.cpp/pull/2053) on a scratch copy: the parent printed `IDENTITY_OK` naming the pinned tree and the child then printed `DECOY ti2vid_one_stage RUNNING -- the parent's assert never saw me`. This is the `.agents/specs/ltx-2-5.md` §7.0(b) decoy failure surviving in the one process that touches weights. The 2026-08-27 reference render is NOT affected and its `gateable = yes` stands: `render.sh` issues no `cd`, `/workspace/ltx2-oracle/` holds no `ltx_*` package, and the committed manifest records module origins inside the pinned clone — so the defect is an OVER-CLAIMED guarantee, not a wrong measurement. Fix: `-P` (or `PYTHONSAFEPATH=1`) on the child plus an explicit `cwd=`, with the reviewer's decoy as the red-first test. NOT fixed in flow and deliberately: the script's sha256 equals the one the worker executed and printed, and that equality is the provenance chain `.agents/oracles/ltx-2.md`'s `gateable = yes` rests on, so editing the file for a hardening that changed no result would trade a verifiable fact for a better comment. Owned by `ENG-UPSTREAM-LTX2-PIN` and listed under `## Owed` in [oracle-ltx-2-pin.md](../specs/oracle-ltx-2-pin.md) | bug | -| [#2056](https://github.com/mudler/vllm.cpp/issues/2056) | `MODEL-MM-QWEN4-EXP` | **`check-agent-record.py` accepts TWO claim files owning the same matrix row**, so a claim collision merges clean and silent. Measured on this branch: copying W6a's `CLAIM-MODEL-MM-QWEN4-EXP.md` beside W1's `CLAIM-MODEL-MM-QWEN4-EXP-W1.md` gives `agent record OK`, rc=0, with both files asserting ownership of `MODEL-MM-qwen4-exp-qwen4-exp-for-conditional-generation` and both marked `ACTIVE`. Git cannot conflict on it because the two sides touch different PATHS. The matrix owner cell holds exactly ONE value, so the record goes silently ambiguous. Resolved here by merge ORDER, which is an operator remembering rather than a gate. NOT fixed in flow: it changes checker semantics and owes its own row, spec and red-before test per AGENTS.md. Listed under `## Owed` in [qwen4-exp-flash-next.md](../specs/qwen4-exp-flash-next.md). | bug | -| [#1973](https://github.com/mudler/vllm.cpp/issues/1973) | `KV-DSV4-MULTICACHE` | **`MakeDeepseekV4KVCache` publishes ONE placeholder group for a model that needs 167 caches, and the runner would drop the real topology in silence.** W2 of [#1925](https://github.com/mudler/vllm.cpp/issues/1925), after W1 ([#1960](https://github.com/mudler/vllm.cpp/issues/1960)) landed the allocation metadata and published none of it. `deepseek_v4_registry.cpp:126-148` emits one `"mla"` group at `head_size = head_dim + qk_rope_head_dim = 576` and calls itself a STUB; upstream at the pin `5559679229bc961848b121ccdeaa8fa5d79bec98` publishes **167 entries in seven groups** across DeepSeek-V4-Flash's 43 layers -- 21 C4A + 20 C128A compressed latents and 21 indexer key caches as `MLAAttentionSpec` (`vllm/models/deepseek_v4/attention.py:631-645`, `:669-684`), 43 SWA caches (`vllm/v1/attention/backends/mla/sparse_swa.py:86-101`) and 21 + 21 + 20 compressor states (`vllm/models/deepseek_v4/compressor.py:188-200`) as `SlidingWindowMLASpec`, at block sizes 256, 64, 4 and 8. **The factory is REACHED from production** -- it is `ModelFactory::make_kv_cache` (`:120`), called from `LoadedEngine`'s constructor through `MakeKVCacheResolved`/`MakeKVCacheMaybeSpec` (`src/vllm/entrypoints/model_loader.cpp:1394-1404`, `:1681`) -- so what it publishes reaches `GPUModelRunner::initialize_kv_cache` unfiltered, and that loop (`src/vllm/v1/worker/gpu/runner.cpp:577-597`) has two arms and no `else`: `kSlidingWindowMla` matches nothing and a SECOND `kMlaAttention` group is passed over by the `full_attn_group_id_ < 0` guard, with no diagnostic either way. **Worse than one dropped group, and this is what W1 could not see from the selection loop alone:** `membership_by_name` is set only inside `if (has_mamba_group)` (`runner.cpp:820-845`), DeepSeek-V4 has no Mamba group, so the allocation loop falls into `is_full_attn = !is_gdn` and allocates ONE buffer per HIDDEN LAYER sized from the target group's page -- 43 buffers of one page for a model needing 167 of seven. A silently short KV allocation is a wrong-tokens failure, not a crash. FIXED by publishing the seven groups WITH a `VT_CHECK` refusal naming any published group the runner does not allocate, which is `AGENTS.md`'s "refuse an unimplemented arm with a message that names the missing part" rather than a new policy. **Consequence stated rather than implied: DeepSeek-V4 on the server path now refuses at engine construction instead of running.** What it loses is a full-recompute path with no decode step -- `Forward`/`ForwardDevice` discard `attn_kv` (`src/vllm/model_executor/models/deepseek_v4.cpp:2886-2887`, `:2959-2960`) -- and `examples/deepseek_v4_gen` does not go through the runner and is unaffected. Two further refusals are NOT in scope and are named so nobody takes them for this: `HybridKVCacheCoordinator`'s uniform-block-size `assert`, inert under `NDEBUG` (`src/vllm/v1/core/kv_cache_coordinator.cpp:340-346`, W4), and the missing third `ModelForwardInput` channel (`include/vllm/model_executor/models/model_registry.h:303-315`, W3). Nothing published is consumed; reachability is owed to W3 and W5 under `## Owed` in [kv-dsv4-multicache.md](../specs/kv-dsv4-multicache.md) | bug | +| [#2059](https://github.com/mudler/vllm.cpp/issues/2059) | — | `.github/workflows/ci.yml:1598` sets `VT_POOL_BYPASS: "1"` for BOTH `sanitize-cpu` lanes, so the `DevicePool` free list, size-class ladder, best-fit borrow (#1922) and capture pre-grow (#1380) are unexecuted under ASan AND TSan. The stated justification is ASan's `detect_leaks`; ThreadSanitizer has no leak detector and gains nothing. MEASURED: under `-DVLLM_CPP_SANITIZE=thread` with the pool ENABLED, `test_qwen3_5_decode_graph_seam` is 10/10, 156 assertions, exit 0, zero TSan warnings. Found while repairing the #2047 red; owed under `## Owed` in [cudagraph-pregrow-nonspec.md](specs/cudagraph-pregrow-nonspec.md) | bug | +| [#2055](https://github.com/mudler/vllm.cpp/issues/2055) | `ENG-UPSTREAM-LTX2-PIN` | `tools/oracle/ltx2_oracle.py` asserts the LTX-2 oracle's revision and its resolved `ltx_core` / `ltx_pipelines` origins in the PARENT, then renders in a child started with `python -m`, which puts the CURRENT WORKING DIRECTORY on that child's `sys.path[0]`. The parent's `importlib.util.find_spec` never consults it, so a directory holding a decoy `ltx_pipelines`, made the CWD, is imported by the process that loads the weights while the process that checked identity sees nothing. MEASURED by a fresh reviewer of [#2053](https://github.com/mudler/vllm.cpp/pull/2053) on a scratch copy: the parent printed `IDENTITY_OK` naming the pinned tree and the child then printed `DECOY ti2vid_one_stage RUNNING -- the parent's assert never saw me`. This is the `.agents/specs/ltx-2-5.md` §7.0(b) decoy failure surviving in the one process that touches weights. The 2026-08-27 reference render is NOT affected and its `gateable = yes` stands: `render.sh` issues no `cd`, `/workspace/ltx2-oracle/` holds no `ltx_*` package, and the committed manifest records module origins inside the pinned clone — so the defect is an OVER-CLAIMED guarantee, not a wrong measurement. Fix: `-P` (or `PYTHONSAFEPATH=1`) on the child plus an explicit `cwd=`, with the reviewer's decoy as the red-first test. NOT fixed in flow and deliberately: the script's sha256 equals the one the worker executed and printed, and that equality is the provenance chain `.agents/oracles/ltx-2.md`'s `gateable = yes` rests on, so editing the file for a hardening that changed no result would trade a verifiable fact for a better comment. Owned by `ENG-UPSTREAM-LTX2-PIN` and listed under `## Owed` in [oracle-ltx-2-pin.md](specs/oracle-ltx-2-pin.md) | bug | +| [#2056](https://github.com/mudler/vllm.cpp/issues/2056) | `MODEL-MM-QWEN4-EXP` | **`check-agent-record.py` accepts TWO claim files owning the same matrix row**, so a claim collision merges clean and silent. Measured on this branch: copying W6a's `CLAIM-MODEL-MM-QWEN4-EXP.md` beside W1's `CLAIM-MODEL-MM-QWEN4-EXP-W1.md` gives `agent record OK`, rc=0, with both files asserting ownership of `MODEL-MM-qwen4-exp-qwen4-exp-for-conditional-generation` and both marked `ACTIVE`. Git cannot conflict on it because the two sides touch different PATHS. The matrix owner cell holds exactly ONE value, so the record goes silently ambiguous. Resolved here by merge ORDER, which is an operator remembering rather than a gate. NOT fixed in flow: it changes checker semantics and owes its own row, spec and red-before test per AGENTS.md. Listed under `## Owed` in [qwen4-exp-flash-next.md](specs/qwen4-exp-flash-next.md). | bug | +| [#1973](https://github.com/mudler/vllm.cpp/issues/1973) | `KV-DSV4-MULTICACHE` | **`MakeDeepseekV4KVCache` publishes ONE placeholder group for a model that needs 167 caches, and the runner would drop the real topology in silence.** W2 of [#1925](https://github.com/mudler/vllm.cpp/issues/1925), after W1 ([#1960](https://github.com/mudler/vllm.cpp/issues/1960)) landed the allocation metadata and published none of it. `deepseek_v4_registry.cpp:126-148` emits one `"mla"` group at `head_size = head_dim + qk_rope_head_dim = 576` and calls itself a STUB; upstream at the pin `5559679229bc961848b121ccdeaa8fa5d79bec98` publishes **167 entries in seven groups** across DeepSeek-V4-Flash's 43 layers -- 21 C4A + 20 C128A compressed latents and 21 indexer key caches as `MLAAttentionSpec` (`vllm/models/deepseek_v4/attention.py:631-645`, `:669-684`), 43 SWA caches (`vllm/v1/attention/backends/mla/sparse_swa.py:86-101`) and 21 + 21 + 20 compressor states (`vllm/models/deepseek_v4/compressor.py:188-200`) as `SlidingWindowMLASpec`, at block sizes 256, 64, 4 and 8. **The factory is REACHED from production** -- it is `ModelFactory::make_kv_cache` (`:120`), called from `LoadedEngine`'s constructor through `MakeKVCacheResolved`/`MakeKVCacheMaybeSpec` (`src/vllm/entrypoints/model_loader.cpp:1394-1404`, `:1681`) -- so what it publishes reaches `GPUModelRunner::initialize_kv_cache` unfiltered, and that loop (`src/vllm/v1/worker/gpu/runner.cpp:577-597`) has two arms and no `else`: `kSlidingWindowMla` matches nothing and a SECOND `kMlaAttention` group is passed over by the `full_attn_group_id_ < 0` guard, with no diagnostic either way. **Worse than one dropped group, and this is what W1 could not see from the selection loop alone:** `membership_by_name` is set only inside `if (has_mamba_group)` (`runner.cpp:820-845`), DeepSeek-V4 has no Mamba group, so the allocation loop falls into `is_full_attn = !is_gdn` and allocates ONE buffer per HIDDEN LAYER sized from the target group's page -- 43 buffers of one page for a model needing 167 of seven. A silently short KV allocation is a wrong-tokens failure, not a crash. FIXED by publishing the seven groups WITH a `VT_CHECK` refusal naming any published group the runner does not allocate, which is `AGENTS.md`'s "refuse an unimplemented arm with a message that names the missing part" rather than a new policy. **Consequence stated rather than implied: DeepSeek-V4 on the server path now refuses at engine construction instead of running.** What it loses is a full-recompute path with no decode step -- `Forward`/`ForwardDevice` discard `attn_kv` (`src/vllm/model_executor/models/deepseek_v4.cpp:2886-2887`, `:2959-2960`) -- and `examples/deepseek_v4_gen` does not go through the runner and is unaffected. Two further refusals are NOT in scope and are named so nobody takes them for this: `HybridKVCacheCoordinator`'s uniform-block-size `assert`, inert under `NDEBUG` (`src/vllm/v1/core/kv_cache_coordinator.cpp:340-346`, W4), and the missing third `ModelForwardInput` channel (`include/vllm/model_executor/models/model_registry.h:303-315`, W3). Nothing published is consumed; reachability is owed to W3 and W5 under `## Owed` in [kv-dsv4-multicache.md](specs/kv-dsv4-multicache.md) | bug | | [#1974](https://github.com/mudler/vllm.cpp/issues/1974) | `KV-DSV4-MULTICACHE` | **`spec_equal`'s `default:` arm returns false for `kMlaAttention` and `kSlidingWindowMla`, so two identical MLA specs never merge into one `SpecGroup`.** Observed by W1 ([#1960](https://github.com/mudler/vllm.cpp/issues/1960)) from a read of the switch and correctly left alone unverified; verified and fixed in flow with W2 ([#1973](https://github.com/mudler/vllm.cpp/issues/1973)). `spec_equal` (`src/vllm/v1/core/kv_cache_coordinator.cpp:17-67`) is our port of the frozen-dataclass `__eq__` and carries arms for `kFullAttention`, `kSlidingWindow`, `kChunkedLocalAttention` and `kMamba` only; `HybridKVCacheCoordinator::verify_and_split_kv_cache_groups` (`:353-376`) uses it to batch groups sharing a spec. Upstream cannot answer `false` there: every spec class is `@dataclass(frozen=True, kw_only=True)` (`vllm/v1/kv_cache_interface.py:380-381`, `:610-611`), so `__eq__` is generated over all fields and two identical `MLAAttentionSpec`s are equal. **Latent rather than observed at runtime, and that is stated rather than implied:** the coordinator needs two groups, and every MLA model in the tree publishes exactly one MLA group (`deepseek_v2_registry.cpp`, `deepseek_v4_registry.cpp`, `glm4_moe_lite_registry.cpp`, `kimi_k3_registry.cpp`, `kimi_linear_registry.cpp`, `minicpm3_registry.cpp`, `dots3_note.cpp`), so the helper is never called on two of them today. It stops being latent with #1973, which publishes three `kMlaAttention` and four `kSlidingWindowMla` groups for DeepSeek-V4. The cost of a wrong `false` is a second `SpecGroup`, so `find_longest_cache_hit` runs per group instead of per distinct spec and the eagle-bit propagation (`:395-402`) is computed over a different partition than upstream's; a wrong `true` would be the dangerous direction, and this is the safe one, which is why it is a mirror divergence rather than a live miscompute. Fixed by adding both arms, comparing `FullAttentionSpec`'s / `SlidingWindowSpec`'s fields plus the four DeepSeek-V4 fields `cache_dtype_str`, `alignment`, `compress_ratio` and `model_version`; `default:` keeps returning `false` for the kinds with no ported struct | bug | | [#1979](https://github.com/mudler/vllm.cpp/issues/1979) | `ENG-RELEASE-WINDOWS` | **`windows-msvc-vulkan` is red on every open PR: `test_openai_api_server.exe` fail-fasts with `0xC0000409`, and the job has no `main` baseline to bisect against.** Found while gating [#1973](https://github.com/mudler/vllm.cpp/issues/1973) (PR [#1977](https://github.com/mudler/vllm.cpp/pull/1977)); PRE-EXISTING and not that row's defect. `scripts/build-windows-release.ps1:31` throws on `build-pr-windows-vulkan\tests\Release\test_openai_api_server.exe exited with status -1073740791`, which is `0xC0000409` `STATUS_STACK_BUFFER_OVERRUN`, the MSVC fail-fast code. It is a CRASH and not a doctest failure: the log carries no `[doctest]` summary and no assertion line for that binary, so nothing reports which case was running; the last lines before it are ordinary chat-completions request logs. MEASURED 2026-08-26 by reading four job logs directly -- [#1977](https://github.com/mudler/vllm.cpp/pull/1977) job 98175568298, [#1975](https://github.com/mudler/vllm.cpp/pull/1975) job 98169254199, [#1969](https://github.com/mudler/vllm.cpp/pull/1969) job 98152969359 and [#1967](https://github.com/mudler/vllm.cpp/pull/1967) job 98144800167 -- all four stopping on the SAME binary. Those four pull requests touch the KV-cache interface, the runner, the LTX-2.5 lane and the record surfaces, and the intersection of their diffs contains no file under `src/vllm/entrypoints/`, so the common factor is the job and not the change. **The second half of the defect is that `windows-msvc-cpu` and `windows-msvc-vulkan` are PR-ONLY jobs that never run on `main`**, so there is no green reference commit to bisect against and no signal for when this started -- which makes a repository-wide red something every author must independently re-attribute, and this issue exists so that stops happening a fifth time. NOT ESTABLISHED and stated rather than implied: no Windows host was available, so nothing was reproduced locally, the crashing case is not identified, the first bad commit is not identified and no cause is proposed; `httplib`'s `stop()` being a no-op on a never-listened server so `SO_REUSEPORT` splits the port (`.agents/environment.md`) is named as an ADJACENT shape to check first, NOT as an attribution. NOT FIXED IN FLOW: it needs a Windows host to reproduce, it sits in a server path unrelated to the KV-cache row that found it, and identifying the crashing case is its own investigation | bug | -| [#1998](https://github.com/mudler/vllm.cpp/issues/1998) | `MODEL-MM-GLM53-FLASH` | **`zai-org/GLM-5.3-Flash` declares `Glm5NextForConditionalGeneration` / `glm5_next`, an architecture vLLM implements at NO revision, whose only admissible oracle is transformers and which NOTHING on this fleet can execute.** Read live 2026-08-26: `git grep "Glm5\|glm5_next"` returns ZERO hits at the parity pin `555967922` AND at vLLM `origin/main` = `c71f6f8a81` (2026-08-26 08:42 -0700), whose `vllm/models/` package holds `common, deepseek_v32, deepseek_v4, dots3_note, inkling, kimi_k3, minimax_m3` and no `glm5next`. That is ABSENCE from vLLM `main`, not staleness in our pin, so a pin advance — 348 commits of reconciliation — reaches nothing here. [vllm#53906](https://github.com/vllm-project/vllm/pull/53906) "[Model] add GLM-5.3-Flash support" (opened 2026-08-26T14:12:00Z, head `933876c388fb`, 85 files, +12,511/-540) would register `Glm5NextForCausalLM`, `Glm5NextForConditionalGeneration` and `Glm5NextMTPModel`, and is **OPEN, unmerged, `mergeable: false`** — an unmerged PR is not a revision and is INADMISSIBLE as an oracle, the same rule this index already applied to SGLang #36497 on the qwen4-exp row; it is cited only as evidence of upstream shape. SGLang implements nothing either (PR #36507 open; code search 0) and the trap there is that its **cookbook** PRs #36440 and #36513 MERGED at 14:00:16Z and 14:39:29Z — documentation landed ahead of the code, so the model card links a recipe SGLang `main` cannot run. llama.cpp implements nothing (code search 0; PR #27752 open; our oracle pinned at `b10451`), and vllm-omni implements nothing. **The sole admissible reference is transformers**, implementing commit `eb4d9e2a64` (2026-08-26T14:26:40Z, transformers#48342), and the first RELEASE carrying it is **`v5.16.1`** — bounded rather than assumed by fetching `models/glm5_next/modeling_glm5_next.py` at each tag: HTTP **200** at `v5.16.1`, **404** at `v5.16.0`, **404** at `v5.15.1`. Our transformers pin is 5.14.1, so this row needs a lane-scoped second pin expiring when vLLM registers `glm5_next`; note it is `5.16.1` where the `MODEL-MM-QWEN4-EXP` lane is `5.16.0`, one day apart because `Qwen4Exp` merged before the 5.16.0 cut and `Glm5Next` after it — two lanes, two releases, which is what a lane pin is for. **Not the blocked GLM-5 row:** `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (`GlmMoeDsaForCausalLM`, `BLOCKED`, `CLAIM-GLM-DSA-LATEST-DEEPSEEK`) is DeepSeek-V3.2 verbatim at 753.9B; `glm5_next` is a different `model_type`, a different parameter count, a different attention topology and has a vision tower. **The architecture:** 45 layers, hidden 4096, vocab 154880, 1,048,576 context, `layer_types` = 34 `linear_attention` (KDA) + 11 `deepseek_sparse_attention` at 3,7,…,43; MoE 288 routed + 1 shared, top-8, `moe_intermediate_size` 2048, sigmoid/`noaux_tc`, `n_group=topk_group=1` (which makes the group stage a NO-OP), `routed_scaling_factor` 2.5, clamped SwiGLU `swiglu_limit` 10.0 in FIVE places; mHC `hc_mult` 4 / `hc_sinkhorn_iters` 20 / `hc_eps` 1e-6; vision 24L/1024/patch **14**/merge 2/temporal 2/out 4096. **The text stack has NO rotary anywhere.** `qk_rope_head_dim: 0`, `mla_use_nope: true`, no `rope_theta`, no `rope_scaling`; the reference *deletes* the inherited `rope_parameters` field and passes `position_embeddings=None` to every layer, and **`indexer_rope_interleave: true` is a VESTIGIAL flag** the indexer override ignores — implementing it because the config names one is implementing a bug. Position comes only from the KDA decay and short convs. **HEAVY REUSE — and the two obvious reuses are WRONG in one place each, in ways a token gate cannot see.** (1) The KDA forget gate takes the OTHER branch: `gate_lower_bound: -5.0` is not None, so it computes `-5.0 * sigmoid(exp(A_log) * (f_b(f_a(x)) + dt_bias))`, whereas our `src/vllm/model_executor/models/kimi_kda.cpp:60` implements Kimi-Linear's `-exp(A_log) * softplus(g + dt_bias)`. Different functions of the same inputs, both smooth, both negative, both fluent — and the sign of `decay_rate` differs too, so getting it wrong flips decay to growth. (2) The mHC head collapse is `hidden_streams.mean(dim=2)`, an UNWEIGHTED mean whose own docstring says "Unlike DeepSeek-V4", where our `src/vllm/model_executor/models/deepseek_v4_mhc.cpp:168` `HcHeadCollapse` is the V4 weighted-gated collapse; the checkpoint carries **no `hc_head.*` tensors at all**, which settles it. **(3) The k-pool indexer is unobservable at or below 2048 tokens:** with `index_topk: 2048` a shorter context selects every candidate, the selection is the identity, the pooling is invisible and any short-prompt gate passes over a broken indexer — the same shape as the trap recorded on the qwen4-exp row. The pooling is **learned, per channel, not a mean**: 128 independent 4-way softmaxes over a learned `[4,128]` intra-pool position embedding plus a cached `[128,4096]` gate score, pools aligned to the first VALID token (so the grid is batch-dependent), a pool must be COMPLETE to be a candidate and is visible iff its LAST member is, `select_k = min(index_topk/index_kpool, P) = min(512, P)`, expanded back to raw token indices and widened to `index_topk + kpool - 1` = **2051** by an unscored ragged tail. Its cache is 257 floats/token/layer, not the parent's 128. **Structural blocker:** `MlaBlockDims::Validate` REFUSES `qk_rope_head_dim == 0` (`src/vllm/model_executor/layers/attention/mla_attention.cpp:90-93` requires every dim > 0, `:95-99` requires it even); Kimi-Linear is NOT the same case — it sets `mla_use_nope` while keeping `qk_rope_head_dim: 64`, so the rope slice still occupies the cache row (`include/vllm/model_executor/models/kimi_linear.h:86,:88`). NoPE does buy one thing: MLA weight absorption becomes trivially valid, since there is no rope slice to keep out of it. **NOTHING FITS, and no GGUF exists.** Measured live by HTTP RANGE over all 62 safetensors headers, payload never fetched: 76,108 tensors, `F8_E4M3` 314,396,639,232 + `BF16` 6,926,096,640 elements = **321,322,735,872 real parameters** plus 19,484,766 F32 block scales; routed experts are **311.65B = 97.0%**, so any GGUF arm's mixed bpw is the experts' bpw to within a percent. On disk against ~119.63 GiB on GB10: `zai-org/GLM-5.3-Flash` FP8 e4m3 block-128x128 **305.78 GiB (2.56x over)**, `zai-org/GLM-5.3-Flash-BF16` **598.53 GiB (5.00x)**, `unsloth/GLM-5.3-Flash-FP8` 305.82 GiB, `LibertAIDAI/GLM-5.3-Flash-NVFP4` **181.32 GiB (1.52x)**. **All four repos named `*-GGUF` contain ZERO `.gguf` files** (`unsloth`, `AtomicChat`, `aj9o9`, `vcruz305` — READMEs and images); a repository name is not an artifact. **`gateable = no` on MEMORY, and the consequence is the scoping result: no oracle can execute this model on any device this project reaches, so NO end-to-end token gate exists or can exist here** — recorded as visible debt, not waived. What replaces it is a tiny-shape CPU reference oracle instantiated from the pinned transformers, which is a real oracle for the NUMERICS and is explicitly NOT oracle gateability for the model. Developer-set W7 target is a low GGUF quant that fits the DGX; arithmetic (not measurement) with experts at the named type and the other 3% at Q6_K, bpw taken from our own reader's block traits at `src/vllm/model_executor/model_loader/gguf_reader.cpp:200`: Q8_0 318.0 GiB, Q6_K 245.5, Q4_K 170.6, IQ3_XXS 118.1 (~1.5 GiB left — not a margin), **Q2_K 102.6 (~17 GiB left)**, IQ2_S 100.3, IQ2_XXS 82.3, IQ1_S 64.0; KV is small by design at ~11.7 KiB/token (1.5 GiB at 128K, ~11.4 GiB at the full 1M) though #1963/#1966 mean that arithmetic is re-derived against the runner, not trusted. **Q2_K is the recommended first arm and the reason is producibility, not quality: i-quants need an importance matrix, an importance matrix needs a forward pass, and a forward pass needs 181 GiB — the dependency is CIRCULAR on this fleet.** K-quants need none. And since llama.cpp has no `glm5_next`, the converter is ours to author, with no reference GGUF to diff against and no llama.cpp floor for the arm. Layer 45 is the MTP block — 888 tensors, 7.43B params, 2.31% of the model, DSA/MLA rather than KDA and carrying no `hc_*` tensors — present in the checkpoint and discarded by the reference at `modular_glm5_next.py:1235`; owed, not implemented, and skipping it in the converter is ~2.3 GiB of headroom. Live seam contention to rebase onto rather than fork: [#1971](https://github.com/mudler/vllm.cpp/pull/1971) (DSA geometry) and [#1977](https://github.com/mudler/vllm.cpp/pull/1977) (DSv4 KV multicache). Spec: [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md), 8 waves, W0-W2 and W4 CPU-gateable and W3/W5-W8 needing a GPU. No product code lands under the spec pull request | feature | -| [#2060](https://github.com/mudler/vllm.cpp/issues/2060) | `MODEL-MM-QWEN4-EXP` | **The `MODEL-MM-QWEN4-EXP` speed gate needs a llama.cpp denominator on the identical artifact, and no released llama.cpp can name the architecture.** Measured 2026-08-27 from refs and objects, never from the developer's dirty checkout: `git grep -il qwen4exp b10451` returns nothing tree-wide at rc=1, with the same grep for `qwen3vl` at rc=0 proving the instrument works. The only llama.cpp that converts or loads `qwen4exp` is ggml-org/llama.cpp [#27742](https://github.com/ggml-org/llama.cpp/pull/27742), still OPEN: `git merge-base --is-ancestor 035e2273 origin/master` is rc=1, and so is the same test on the live head. The `llama-cpp` file cannot carry it, because `check-oracle-pins.py` admits one `oracle-pin` block per file and that file's pin is deliberately stock release `b10451`, the floor being what a user can get today. **This issue was STALE AT BIRTH, which is the exact reason a PR needs a recorded head SHA rather than a branch name**: it names `035e22731a7fd70b9854b3a2d64ec68e9b1a45d3` as the head, and the head had already been `6c5afc86ae84448ae4d744e357017e2c490ad9c3` for 3 h 55 m when the issue was created. Measured from the forge: issue created `2026-08-27T07:53:09Z`, `6c5afc86` committed `2026-08-27T03:58:12Z`, `035e2273` committed `2026-08-26T15:09:53Z`. Nothing drifted while the issue sat, and the 12 h 48 m is the span between the two commits. The gap is a clean fast-forward with nothing force-pushed away. Landed as `.agents/oracles/llama-cpp-qwen4exp.md` at `035e2273` because that is the revision the build evidence was measured at, and a pin whose evidence was measured at another revision is not a pin. Advancing to `6c5afc86` breaks NO anchor in this tree and would GAIN ONE `-Werror` build fix, `6a69a0c` (NOT two: `24ea62d` says "and the fatal-warning build" in its subject, but measured from the object in a fresh bare clone its diff touches one file, `src/llama-memory-hybrid-idx.cpp`, repairing a `ple_hist_rm` iterator invalidation and tightening an `n_toks` bound, and the `mem_size` line had been gone for three commits by then, `grep -c mem_size` returning 1 at `6a69a0c~1` and 0 at `6a69a0c`): the spec's three `conversion/qwen.py` anchors (`:365`, `:387-388`, `:438`) resolve at stock `b10451`, which [qwen4-exp-flash-next.md](../specs/qwen4-exp-flash-next.md) states itself; `conversion/qwen.py` is byte-identical across the two revisions, blob `cdba8a63e9c919232e2ec80e88b01afec7967dc4` at both; the two facts that spec does read at the PR, `conversion/qwen4exp.py:19` declaring `class Qwen4ExpTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase)` and a `modify_tensors` with no `hc_norm` branch, hold identically at `6c5afc86`; and `fbe1773` and `5674c73`, the commits that touch tensor naming and the QKV layout, touch `src/` files only, which no anchor cites. The hold keeps the pin and its evidence pointing at one object. It is not a claim that advancing is unsafe. **`gateable = no`, and this issue owes the run half.** Both halves were attempted rather than asserted: a CPU-only build from a fresh `git archive` of the pinned object SUCCEEDS, 247 translation units, `libllama.so` carrying 228 `qwen4exp` strings so the architecture links rather than merely existing in source, evidence in [docs/bench-evidence/oracle-llamacpp-qwen4exp-pr27742-build-20260827.md](../../docs/bench-evidence/oracle-llamacpp-qwen4exp-pr27742-build-20260827.md), and the one warning it emits is the unused `mem_size` that a later commit removes as an `-Werror` break, which is the gap between the pin and the head made concrete. The RUN half is blocked on artifact availability and not on the oracle: the seven `unsloth/Qwen3.8-Flash-Next-GGUF` quants were mid-download, `UD-IQ1_S` held 3.3 GiB of 67.56 GiB with a 10 MiB fragment of shard 1 of 3, and the attempted load refuses by naming the missing shard 2, which says nothing about whether #27742 loads the architecture. A build proves the architecture is declared, never that its graph produces coherent text | record | +| [#1998](https://github.com/mudler/vllm.cpp/issues/1998) | `MODEL-MM-GLM53-FLASH` | **`zai-org/GLM-5.3-Flash` declares `Glm5NextForConditionalGeneration` / `glm5_next`, an architecture vLLM implements at NO revision, whose only admissible oracle is transformers and which NOTHING on this fleet can execute.** Read live 2026-08-26: `git grep "Glm5\|glm5_next"` returns ZERO hits at the parity pin `555967922` AND at vLLM `origin/main` = `c71f6f8a81` (2026-08-26 08:42 -0700), whose `vllm/models/` package holds `common, deepseek_v32, deepseek_v4, dots3_note, inkling, kimi_k3, minimax_m3` and no `glm5next`. That is ABSENCE from vLLM `main`, not staleness in our pin, so a pin advance — 348 commits of reconciliation — reaches nothing here. [vllm#53906](https://github.com/vllm-project/vllm/pull/53906) "[Model] add GLM-5.3-Flash support" (opened 2026-08-26T14:12:00Z, head `933876c388fb`, 85 files, +12,511/-540) would register `Glm5NextForCausalLM`, `Glm5NextForConditionalGeneration` and `Glm5NextMTPModel`, and is **OPEN, unmerged, `mergeable: false`** — an unmerged PR is not a revision and is INADMISSIBLE as an oracle, the same rule this index already applied to SGLang #36497 on the qwen4-exp row; it is cited only as evidence of upstream shape. SGLang implements nothing either (PR #36507 open; code search 0) and the trap there is that its **cookbook** PRs #36440 and #36513 MERGED at 14:00:16Z and 14:39:29Z — documentation landed ahead of the code, so the model card links a recipe SGLang `main` cannot run. llama.cpp implements nothing (code search 0; PR #27752 open; our oracle pinned at `b10451`), and vllm-omni implements nothing. **The sole admissible reference is transformers**, implementing commit `eb4d9e2a64` (2026-08-26T14:26:40Z, transformers#48342), and the first RELEASE carrying it is **`v5.16.1`** — bounded rather than assumed by fetching `models/glm5_next/modeling_glm5_next.py` at each tag: HTTP **200** at `v5.16.1`, **404** at `v5.16.0`, **404** at `v5.15.1`. Our transformers pin is 5.14.1, so this row needs a lane-scoped second pin expiring when vLLM registers `glm5_next`; note it is `5.16.1` where the `MODEL-MM-QWEN4-EXP` lane is `5.16.0`, one day apart because `Qwen4Exp` merged before the 5.16.0 cut and `Glm5Next` after it — two lanes, two releases, which is what a lane pin is for. **Not the blocked GLM-5 row:** `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (`GlmMoeDsaForCausalLM`, `BLOCKED`, `CLAIM-GLM-DSA-LATEST-DEEPSEEK`) is DeepSeek-V3.2 verbatim at 753.9B; `glm5_next` is a different `model_type`, a different parameter count, a different attention topology and has a vision tower. **The architecture:** 45 layers, hidden 4096, vocab 154880, 1,048,576 context, `layer_types` = 34 `linear_attention` (KDA) + 11 `deepseek_sparse_attention` at 3,7,…,43; MoE 288 routed + 1 shared, top-8, `moe_intermediate_size` 2048, sigmoid/`noaux_tc`, `n_group=topk_group=1` (which makes the group stage a NO-OP), `routed_scaling_factor` 2.5, clamped SwiGLU `swiglu_limit` 10.0 in FIVE places; mHC `hc_mult` 4 / `hc_sinkhorn_iters` 20 / `hc_eps` 1e-6; vision 24L/1024/patch **14**/merge 2/temporal 2/out 4096. **The text stack has NO rotary anywhere.** `qk_rope_head_dim: 0`, `mla_use_nope: true`, no `rope_theta`, no `rope_scaling`; the reference *deletes* the inherited `rope_parameters` field and passes `position_embeddings=None` to every layer, and **`indexer_rope_interleave: true` is a VESTIGIAL flag** the indexer override ignores — implementing it because the config names one is implementing a bug. Position comes only from the KDA decay and short convs. **HEAVY REUSE — and the two obvious reuses are WRONG in one place each, in ways a token gate cannot see.** (1) The KDA forget gate takes the OTHER branch: `gate_lower_bound: -5.0` is not None, so it computes `-5.0 * sigmoid(exp(A_log) * (f_b(f_a(x)) + dt_bias))`, whereas our `src/vllm/model_executor/models/kimi_kda.cpp:60` implements Kimi-Linear's `-exp(A_log) * softplus(g + dt_bias)`. Different functions of the same inputs, both smooth, both negative, both fluent — and the sign of `decay_rate` differs too, so getting it wrong flips decay to growth. (2) The mHC head collapse is `hidden_streams.mean(dim=2)`, an UNWEIGHTED mean whose own docstring says "Unlike DeepSeek-V4", where our `src/vllm/model_executor/models/deepseek_v4_mhc.cpp:168` `HcHeadCollapse` is the V4 weighted-gated collapse; the checkpoint carries **no `hc_head.*` tensors at all**, which settles it. **(3) The k-pool indexer is unobservable at or below 2048 tokens:** with `index_topk: 2048` a shorter context selects every candidate, the selection is the identity, the pooling is invisible and any short-prompt gate passes over a broken indexer — the same shape as the trap recorded on the qwen4-exp row. The pooling is **learned, per channel, not a mean**: 128 independent 4-way softmaxes over a learned `[4,128]` intra-pool position embedding plus a cached `[128,4096]` gate score, pools aligned to the first VALID token (so the grid is batch-dependent), a pool must be COMPLETE to be a candidate and is visible iff its LAST member is, `select_k = min(index_topk/index_kpool, P) = min(512, P)`, expanded back to raw token indices and widened to `index_topk + kpool - 1` = **2051** by an unscored ragged tail. Its cache is 257 floats/token/layer, not the parent's 128. **Structural blocker:** `MlaBlockDims::Validate` REFUSES `qk_rope_head_dim == 0` (`src/vllm/model_executor/layers/attention/mla_attention.cpp:90-93` requires every dim > 0, `:95-99` requires it even); Kimi-Linear is NOT the same case — it sets `mla_use_nope` while keeping `qk_rope_head_dim: 64`, so the rope slice still occupies the cache row (`include/vllm/model_executor/models/kimi_linear.h:86,:88`). NoPE does buy one thing: MLA weight absorption becomes trivially valid, since there is no rope slice to keep out of it. **NOTHING FITS, and no GGUF exists.** Measured live by HTTP RANGE over all 62 safetensors headers, payload never fetched: 76,108 tensors, `F8_E4M3` 314,396,639,232 + `BF16` 6,926,096,640 elements = **321,322,735,872 real parameters** plus 19,484,766 F32 block scales; routed experts are **311.65B = 97.0%**, so any GGUF arm's mixed bpw is the experts' bpw to within a percent. On disk against ~119.63 GiB on GB10: `zai-org/GLM-5.3-Flash` FP8 e4m3 block-128x128 **305.78 GiB (2.56x over)**, `zai-org/GLM-5.3-Flash-BF16` **598.53 GiB (5.00x)**, `unsloth/GLM-5.3-Flash-FP8` 305.82 GiB, `LibertAIDAI/GLM-5.3-Flash-NVFP4` **181.32 GiB (1.52x)**. **All four repos named `*-GGUF` contain ZERO `.gguf` files** (`unsloth`, `AtomicChat`, `aj9o9`, `vcruz305` — READMEs and images); a repository name is not an artifact. **`gateable = no` on MEMORY, and the consequence is the scoping result: no oracle can execute this model on any device this project reaches, so NO end-to-end token gate exists or can exist here** — recorded as visible debt, not waived. What replaces it is a tiny-shape CPU reference oracle instantiated from the pinned transformers, which is a real oracle for the NUMERICS and is explicitly NOT oracle gateability for the model. Developer-set W7 target is a low GGUF quant that fits the DGX; arithmetic (not measurement) with experts at the named type and the other 3% at Q6_K, bpw taken from our own reader's block traits at `src/vllm/model_executor/model_loader/gguf_reader.cpp:200`: Q8_0 318.0 GiB, Q6_K 245.5, Q4_K 170.6, IQ3_XXS 118.1 (~1.5 GiB left — not a margin), **Q2_K 102.6 (~17 GiB left)**, IQ2_S 100.3, IQ2_XXS 82.3, IQ1_S 64.0; KV is small by design at ~11.7 KiB/token (1.5 GiB at 128K, ~11.4 GiB at the full 1M) though #1963/#1966 mean that arithmetic is re-derived against the runner, not trusted. **Q2_K is the recommended first arm and the reason is producibility, not quality: i-quants need an importance matrix, an importance matrix needs a forward pass, and a forward pass needs 181 GiB — the dependency is CIRCULAR on this fleet.** K-quants need none. And since llama.cpp has no `glm5_next`, the converter is ours to author, with no reference GGUF to diff against and no llama.cpp floor for the arm. Layer 45 is the MTP block — 888 tensors, 7.43B params, 2.31% of the model, DSA/MLA rather than KDA and carrying no `hc_*` tensors — present in the checkpoint and discarded by the reference at `modular_glm5_next.py:1235`; owed, not implemented, and skipping it in the converter is ~2.3 GiB of headroom. Live seam contention to rebase onto rather than fork: [#1971](https://github.com/mudler/vllm.cpp/pull/1971) (DSA geometry) and [#1977](https://github.com/mudler/vllm.cpp/pull/1977) (DSv4 KV multicache). Spec: [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md), 8 waves, W0-W2 and W4 CPU-gateable and W3/W5-W8 needing a GPU. No product code lands under the spec pull request | feature | +| [#2060](https://github.com/mudler/vllm.cpp/issues/2060) | `MODEL-MM-QWEN4-EXP` | **The `MODEL-MM-QWEN4-EXP` speed gate needs a llama.cpp denominator on the identical artifact, and no released llama.cpp can name the architecture.** Measured 2026-08-27 from refs and objects, never from the developer's dirty checkout: `git grep -il qwen4exp b10451` returns nothing tree-wide at rc=1, with the same grep for `qwen3vl` at rc=0 proving the instrument works. The only llama.cpp that converts or loads `qwen4exp` is ggml-org/llama.cpp [#27742](https://github.com/ggml-org/llama.cpp/pull/27742), still OPEN: `git merge-base --is-ancestor 035e2273 origin/master` is rc=1, and so is the same test on the live head. The `llama-cpp` file cannot carry it, because `check-oracle-pins.py` admits one `oracle-pin` block per file and that file's pin is deliberately stock release `b10451`, the floor being what a user can get today. **This issue was STALE AT BIRTH, which is the exact reason a PR needs a recorded head SHA rather than a branch name**: it names `035e22731a7fd70b9854b3a2d64ec68e9b1a45d3` as the head, and the head had already been `6c5afc86ae84448ae4d744e357017e2c490ad9c3` for 3 h 55 m when the issue was created. Measured from the forge: issue created `2026-08-27T07:53:09Z`, `6c5afc86` committed `2026-08-27T03:58:12Z`, `035e2273` committed `2026-08-26T15:09:53Z`. Nothing drifted while the issue sat, and the 12 h 48 m is the span between the two commits. The gap is a clean fast-forward with nothing force-pushed away. Landed as `.agents/oracles/llama-cpp-qwen4exp.md` at `035e2273` because that is the revision the build evidence was measured at, and a pin whose evidence was measured at another revision is not a pin. Advancing to `6c5afc86` breaks NO anchor in this tree and would GAIN ONE `-Werror` build fix, `6a69a0c` (NOT two: `24ea62d` says "and the fatal-warning build" in its subject, but measured from the object in a fresh bare clone its diff touches one file, `src/llama-memory-hybrid-idx.cpp`, repairing a `ple_hist_rm` iterator invalidation and tightening an `n_toks` bound, and the `mem_size` line had been gone for three commits by then, `grep -c mem_size` returning 1 at `6a69a0c~1` and 0 at `6a69a0c`): the spec's three `conversion/qwen.py` anchors (`:365`, `:387-388`, `:438`) resolve at stock `b10451`, which [qwen4-exp-flash-next.md](specs/qwen4-exp-flash-next.md) states itself; `conversion/qwen.py` is byte-identical across the two revisions, blob `cdba8a63e9c919232e2ec80e88b01afec7967dc4` at both; the two facts that spec does read at the PR, `conversion/qwen4exp.py:19` declaring `class Qwen4ExpTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase)` and a `modify_tensors` with no `hc_norm` branch, hold identically at `6c5afc86`; and `fbe1773` and `5674c73`, the commits that touch tensor naming and the QKV layout, touch `src/` files only, which no anchor cites. The hold keeps the pin and its evidence pointing at one object. It is not a claim that advancing is unsafe. **`gateable = no`, and this issue owes the run half.** Both halves were attempted rather than asserted: a CPU-only build from a fresh `git archive` of the pinned object SUCCEEDS, 247 translation units, `libllama.so` carrying 228 `qwen4exp` strings so the architecture links rather than merely existing in source, evidence in [docs/bench-evidence/oracle-llamacpp-qwen4exp-pr27742-build-20260827.md](../docs/bench-evidence/oracle-llamacpp-qwen4exp-pr27742-build-20260827.md), and the one warning it emits is the unused `mem_size` that a later commit removes as an `-Werror` break, which is the gap between the pin and the head made concrete. The RUN half is blocked on artifact availability and not on the oracle: the seven `unsloth/Qwen3.8-Flash-Next-GGUF` quants were mid-download, `UD-IQ1_S` held 3.3 GiB of 67.56 GiB with a 10 MiB fragment of shard 1 of 3, and the attempted load refuses by naming the missing shard 2, which says nothing about whether #27742 loads the architecture. A build proves the architecture is declared, never that its graph produces coherent text | record | | [#2063](https://github.com/mudler/vllm.cpp/issues/2063) | `MODEL-MM-QWEN4-EXP` | **`scripts/audit-live-rows.py` classifies an `ACTIVE` row from `git log` alone, so ANY commit that merely NAMES a row ID changes that row's verdict, and an unrelated row's commit body has already flipped one twice in one day.** The row is the `qwen4_exp` matrix row at `.agents/model-matrix.md:514` (campaign row `MODEL-MM-QWEN4-EXP`, [#1978](https://github.com/mudler/vllm.cpp/issues/1978)). MEASURED on detached worktrees with nothing applied. At `origin/main` `327211c9c` it reads `ABANDONED`, reason `no branch and no commit on main mentioning the row ID`, `scripts/audit-live-rows.py --check` exits 1 on `259 live rows; 1 abandoned ACTIVE`, and `test_shipped_record_has_no_abandoned_active_row` fails on exactly that ID, 57 tests / 1 failure, so `scripts/agent-ready.py` reported both `audit-live-rows` and `test_audit_live_rows` FAILED on every branch merging that main. At `origin/main` `c714b0234`, two commits later, the SAME row reads `LANDED` and the gate is GREEN. Nothing about the row changed. The cure is that `c714b0234`, a `MODEL-MM-GLM53-FLASH` W7a commit, happens to name the qwen4-exp row ID once in its body. **Both verdicts are wrong and the second is the dangerous one**: the row is `ACTIVE` with waves in flight, and `classify_active`'s own comment says a live claim reported as finished is the false negative it exists to prevent. The underlying defect is that this campaign's commits name the CAMPAIGN row `MODEL-MM-QWEN4-EXP` and never the matrix row, so the probe has no true signal and reads whatever noise reaches it: `git log --grep` for the matrix ID over `327211c9c` is EMPTY while `--grep 'MODEL-MM-QWEN4-EXP'` returns W1, W4 and W6a. A lifecycle move therefore does NOT fix it, which separates this from the [#1469](https://github.com/mudler/vllm.cpp/issues/1469) class (#787, #1469, #1504), where the state was genuinely stale. Two consequences for anyone reading this gate: a GREEN reading can measure nothing, and because `squash_merge_commit_message = PR_BODY` a pull request body that quotes a row ID silently changes that row's verdict on `main`. The pull request that filed this issue keeps the matrix ID out of its commit messages and its body for that reason, which is why this row spells it as a path and a campaign name instead. NOT FIXED IN FLOW: found while running `scripts/agent-ready.py` for `MODEL-MM-dots3-note` W4b-3a, a record-only unit with no authority over the Qwen4-Exp campaign, and a checker-semantics change owes its own row, spec, red-before test and green-after evidence. Owned by row `MODEL-MM-QWEN4-EXP` | bug | | [#2061](https://github.com/mudler/vllm.cpp/issues/2061) | `MODEL-TEXT-laguna-laguna-for-causal-lm` | Laguna's grouped MoE issues two `LqGemmGrouped` calls over the same activation, so it quantizes to Q8_K twice where `vt::MoeGateUpSwiGLUGrouped` quantizes once — W11 measured `QuantizeQ8KKernel` at 12.4% of decode GPU. Also a shared-seam obligation, since AGENTS.md routes mergeable MLP projections through the fused group. Bounded by whether a DYNAMIC UD quant gives both expert towers the same block-quant dtype, which W1 measures before any code | perf | -| [#2011](https://github.com/mudler/vllm.cpp/issues/2011) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash has no GGUF, no upstream tool can make one, and upstream Python cannot even quantize: author the `glm5_next` converter and gate its k-quant encoders BYTE-FOR-BYTE against llama.cpp `b10451`.** W7a of [#1998](https://github.com/mudler/vllm.cpp/issues/1998); spec [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) §W7a. Every GPU gate on this row is blocked behind an artifact that does not exist: measured live 2026-08-26 against ~119.63 GiB usable on `dgx:gpu0`, the published arms are FP8 **305.78 GiB**, BF16 **598.53 GiB** and `LibertAIDAI/...-NVFP4` **181.32 GiB**, and all four repositories named `*-GGUF` contain **ZERO `.gguf` files**. **D6 holds but was too strong, and the correction is what made this wave tractable.** llama.cpp implements no `glm5_next` — re-verified at `origin/master` `539f24529` (fetched 2026-08-26) and at our pin: the enumerators are `LLM_ARCH_GLM4`, `LLM_ARCH_GLM4_MOE`, `LLM_ARCH_GLM_DSA` (`src/llama-arch.h:86-88`) and `src/models/glm-dsa.cpp` is **GLM-5.2**, citing `zai-org/GLM-5.2/blob/main/config.json`, a different model. **But every GGUF convention this converter needs is already present AT `b10451`, so no pin advance was taken and none is owed:** `gguf-py/gguf/constants.py:262-264` carries `class KDA` with `{arch}.kda.head_dim` and **`{arch}.kda.gate_lower_bound`** (`KDA.SAFE_GATE` is `master`-only and GLM-5.3-Flash declares no `safe_gate`, so nothing reaches for it); the KDA tensor spellings `ssm_conv1d_q/k/v`, `ssm_f_a/f_b`, `ssm_g_a/g_b`, `ssm_beta`, `ssm_a`, `ssm_dt`, `ssm_norm` are at `src/llama-arch.cpp:465-479`; `gguf-py/gguf/tensor_mapping.py:896-933` maps them from Kimi-Linear's HF module paths, which are **GLM-5.3-Flash's paths verbatim**; and the indexer names, including `indexer_compressor_ape`/`indexer_compressor_gate` for the k-pool stage, are at `:626-636`. **The one thing upstream cannot supply at any revision is an ENCODER:** `gguf.quants.Q2_K` implements `dequantize_blocks` and **no** `quantize_blocks`, so Q2_K/Q6_K/Q8_0 were ported from `ggml/src/ggml-quants.c` @ `b10451` (`:891`, `:1869`, `:276`, over `make_qkx2_quants:799`, `make_qx_quants:628`, `nearest_int:621`) and gated **byte-identical** against a frozen golden captured from that reference compiled `-ffp-contract=off`. Two traps changed bytes and are recorded in the source rather than left to be re-found: `nearest_int` is the `+12582912.0` add-and-mask trick and rounds half to EVEN, not `round`; and C `roundf` in `quantize_row_q8_0_ref` rounds half AWAY FROM ZERO where `np.rint` rounds half to even, so `np.rint` mis-encodes every exact `.5` — caught by a crafted `[0.5, -0.5, 1.5, ...]` case, `[1,-1,2,-2,3,-3]` expected against `[0,0,2,-2,2,-2]` observed. **Exact arithmetic, from the converter's OWN type resolver run over the real topology** rather than bpw times a parameter count: 1719 output tensors carrying **313,890,512,702** parameters — 321.32B less the 7.43B MTP block, which independently confirms the skip is the 2.31% the spec measured. Q2_K (experts Q2_K, everything else Q6_K) **100.35 GiB** at 2.746 mixed bpw, Q6_K 239.89, Q8_0 310.67, BF16 584.67. Against 119.63 GiB at 128K context and one sequence: KV **1.43 GiB** (11 MLA layers x `kv_lora_rank` 512 x 2 B plus an 11 x 64 B indexer side cache = 11,968 B/token), KDA recurrent state **0.14 GiB** (64 x 128 x 128 x 4 B x 34) plus conv states, leaving **~17.7 GiB** — arithmetic, not measurement, and [#1963](https://github.com/mudler/vllm.cpp/issues/1963)/[#1966](https://github.com/mudler/vllm.cpp/issues/1966) mean W5 re-derives it from the runner. The Q2_K figure is **100.35 and not the spec's 102.6** because that table stated every arm including layer 45. The fallback ladder is `want -> Q8_0 -> F32` rather than `want -> F32`, because with an F32 fallback a FINER arm can come out LARGER than a coarser one and a size table must not have that property; on the real model the F32 sliver is 0.08 GiB (the depthwise conv kernels at `ne0=4`, the 4-D downsample and 5-D patch-embed). The tensor inventory is EXACT, not inferred: the real 8.4 MB `model.safetensors.index.json` (76,108 entries) and shard headers 2/32/62 were read by HTTP RANGE, payload never fetched, confirming the three separate `{q,k,v}_conv1d` convs, `hc_{attn,ffn}_{fn,base,scale}` flat on the layer with **no `hc_head.*` at any layer**, `indexer.k_norm.bias` present (settling LayerNorm-with-bias over RMSNorm), and `index_kpool_compress_{ape,gate}` on 12 layers. Refusals by name for every unimplemented arm: the i-quants because an imatrix needs a forward pass, a forward pass needs 181 GiB, and the dependency is **circular on this fleet** (O5); Q3_K/Q4_K/Q5_K because their encoders are unported and shipping an ungated encoder is worse than refusing (O8); `--keep-mtp` because nothing on the loader side reads an MTP tail (O2); a non-`glm5_next` config; and an FP8 tensor whose `weight_scale_inv` companion is missing, which would otherwise produce a loadable, wrong file. Gate: 50 assertions on a SYNTHETIC tiny-shape checkpoint with an independent in-test GGUF reader — no real checkpoint, no GPU, no C++ build. **NOT REACHED and disclosed:** `glm5next` has no `general.architecture` dispatch entry (`src/vllm/entrypoints/model_loader.cpp:1000`), so the file the converter writes is not loadable by this tree; W1 owns that wiring, tracked as **O9** under #1998. **Owed on close:** O7, artifact production — needs the 300-600 GiB checkpoint staged on local disk, explicit developer authority for the download, and a box with room for source and output at once, plus the sha256, recipe and peak RSS the spec's §Evidence requires | feature | -| [#2068](https://github.com/mudler/vllm.cpp/issues/2068) | `KV-DSV4-MULTICACHE` | **`GPUModelRunner` carries at most ONE attention group and ONE cache per layer, so the seven-group topology W2 publishes is refused instead of allocated, and `ModelForwardInput` has no third channel that could carry it.** W3 of [#1925](https://github.com/mudler/vllm.cpp/issues/1925), after W1 ([#1960](https://github.com/mudler/vllm.cpp/issues/1960)) landed `SlidingWindowMLASpec` and the four `MLAAttentionSpec` fields and W2 ([#1973](https://github.com/mudler/vllm.cpp/issues/1973)) made `MakeDeepseekV4KVCache` publish DeepSeek-V4-Flash's real **167 entries in seven groups**; spec [`specs/kv-dsv4-multicache.md`](../specs/kv-dsv4-multicache.md) `### W3 design`. **Three parts, one defect.** (1) Group selection has two arms and no `else` (`src/vllm/v1/worker/gpu/runner.cpp:608-627`): `full_attn_group_id_` takes the FIRST non-eagle `kFullAttention`/`kMlaAttention` group, `gdn_group_id_` a `kMamba` group, and `kSlidingWindowMla`, `kSlidingWindow`, `kChunkedLocalAttention` and a SECOND `kMlaAttention` group match nothing; both ids are plain `int`s (`include/vllm/v1/worker/gpu/runner.h:591-592`). (2) The allocation loop allocates one buffer per HIDDEN LAYER, because `membership_by_name` is computed only inside `if (has_mamba_group)` (`runner.cpp:942-957`) and DeepSeek-V4 has no Mamba group, so it falls into `is_full_attn = !is_gdn`. MEASURED by W2 on a probe binary with the refusal disabled: two groups published, `full_attn_group_id = 0`, `gdn_group_id = -1`, `attn_kv().size() = 4` for 4 hidden layers, and the dropped `kSlidingWindowMla` group's own `page_size_bytes() = 37440` received nothing, with no exception and no message. The per-layer taxonomy that records the routing is a three-valued enum (`runner.h:379-383`) and cannot say `this layer has four caches`. (3) `ModelForwardInput` carries exactly two cache references (`include/vllm/model_executor/models/model_registry.h:344-345`); `attn_kv`'s positional convention cannot address a cache set keyed by (layer x role) the way upstream's `compilation_config.static_forward_context` is keyed by layer NAME. **Upstream has no equivalent generalization to port:** its runner allocates per registered NAME from the start (`vllm/v1/worker/gpu_model_runner.py:7785-7801` walks every `AttentionLayerBase`), so one C4A layer contributes four keys and there is no one-cache-per-layer assumption to relax; what is mirrored is the KEY, which is why the third channel carries the `static_forward_context` name beside each cache. Pin `5559679229bc961848b121ccdeaa8fa5d79bec98`. **FIXED** by entering a generalized path only when the published topology leaves groups over after the target attention group, the recurrent group and the `fa_draft` slot -- exactly the set W2's refusal already computes -- so every model shipping today runs the legacy path untouched, on the `per_layer_attn_specs` empty-means-unchanged contract (`include/vllm/v1/kv_cache_interface.h:538-556`). **Byte-neutrality re-derived by sweep at base `c714b0234` rather than inherited: 36 group-emplacement sites across 31 factory files -- 25 `FullAttentionSpec`, 7 `MLAAttentionSpec`, 3 `MambaSpec`, 1 `SlidingWindowMLASpec`**; by runtime shape 27 single-group factory FILES, `kimi_linear` (MLA+Mamba), `nemotron_h` (FA+Mamba), `qwen3_5_common` (FA+Mamba+`fa_draft` behind `if (num_spec > 0)`) and `deepseek_v4`. The count differs from W2's `34 across 32` because W2 itself added the two DeepSeek-V4 sites. **FILES, not architectures:** several single-group files back more than one `REGISTER_VLLM_MODEL` (`gemma4`/`gemma4_unified`, `olmo2`/`olmo3`, `llama_dense`/`internlm3_llama`, `muse_glimmer`/`muse_glimmer_mm`, the three parakeets, and `llama_model_embedding` reusing `MakeLlamaForCausalLMKVCache`), so counted by REGISTERED ARCHITECTURE it is **42 total = 34 single-group + 7 multi-group + 1 that publishes nothing** (`qwen4_exp`, whose spec factory throws by name). Only `deepseek_v4` has leftovers -- and **seven is its MAXIMUM, not its count**: `add_mla`/`add_swa_mla` return early on empty names, so the published count is a function of the checkpoint's `compress_ratios` and ranges over **1..7**; an all-`ratio == 1` config publishes ONE group, the SWA group, which is a `SlidingWindowMLASpec` that `full_attn_group_id_` never binds to and is therefore itself the leftover. **Every reachable DeepSeek-V4 config is multi-cache.** The full gate includes the SACRED `test_qwen35_paged_engine` regression. **Consequence stated rather than implied: a DeepSeek-V4 engine at `--block-size` 128 or 256 now CONSTRUCTS and allocates all 167 buffers, and its FIRST forward refuses**, because `ModelRegistry::Forward` will not let a multi-cache topology be silently ignored by a forward that discards `attn_kv`. That refusal is what keeps the third channel from landing dead: it reads the channel's payload -- cache count, group count, first layer name -- not its nullness. **W2's refusal is KEPT and still reaches four shapes:** a group whose layer names do not all resolve, a second `kMamba` group, a group whose spec is neither an `AttentionSpec` nor a `MambaSpec`, and an EAGLE group ([#2084](https://github.com/mudler/vllm.cpp/issues/2084), a review repair -- the first cut enumerated three and tested three, and an eagle group therefore passed the refusal and then got no buffer). **Not in scope, named so nobody takes them for this:** the uniform-`block_size` deferral in `HybridKVCacheCoordinator` (`src/vllm/v1/core/kv_cache_coordinator.cpp:340-346`, an `assert` and therefore inert under `NDEBUG`), the block-table geometry and the `KVBytesPerBlock` / `recurrent_state_bytes` budget helpers that still count one page per hidden layer, all W4; and the forward consuming the caches, W5. Both listed under `## Owed` | bug | +| [#2011](https://github.com/mudler/vllm.cpp/issues/2011) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash has no GGUF, no upstream tool can make one, and upstream Python cannot even quantize: author the `glm5_next` converter and gate its k-quant encoders BYTE-FOR-BYTE against llama.cpp `b10451`.** W7a of [#1998](https://github.com/mudler/vllm.cpp/issues/1998); spec [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) §W7a. Every GPU gate on this row is blocked behind an artifact that does not exist: measured live 2026-08-26 against ~119.63 GiB usable on `dgx:gpu0`, the published arms are FP8 **305.78 GiB**, BF16 **598.53 GiB** and `LibertAIDAI/...-NVFP4` **181.32 GiB**, and all four repositories named `*-GGUF` contain **ZERO `.gguf` files**. **D6 holds but was too strong, and the correction is what made this wave tractable.** llama.cpp implements no `glm5_next` — re-verified at `origin/master` `539f24529` (fetched 2026-08-26) and at our pin: the enumerators are `LLM_ARCH_GLM4`, `LLM_ARCH_GLM4_MOE`, `LLM_ARCH_GLM_DSA` (`src/llama-arch.h:86-88`) and `src/models/glm-dsa.cpp` is **GLM-5.2**, citing `zai-org/GLM-5.2/blob/main/config.json`, a different model. **But every GGUF convention this converter needs is already present AT `b10451`, so no pin advance was taken and none is owed:** `gguf-py/gguf/constants.py:262-264` carries `class KDA` with `{arch}.kda.head_dim` and **`{arch}.kda.gate_lower_bound`** (`KDA.SAFE_GATE` is `master`-only and GLM-5.3-Flash declares no `safe_gate`, so nothing reaches for it); the KDA tensor spellings `ssm_conv1d_q/k/v`, `ssm_f_a/f_b`, `ssm_g_a/g_b`, `ssm_beta`, `ssm_a`, `ssm_dt`, `ssm_norm` are at `src/llama-arch.cpp:465-479`; `gguf-py/gguf/tensor_mapping.py:896-933` maps them from Kimi-Linear's HF module paths, which are **GLM-5.3-Flash's paths verbatim**; and the indexer names, including `indexer_compressor_ape`/`indexer_compressor_gate` for the k-pool stage, are at `:626-636`. **The one thing upstream cannot supply at any revision is an ENCODER:** `gguf.quants.Q2_K` implements `dequantize_blocks` and **no** `quantize_blocks`, so Q2_K/Q6_K/Q8_0 were ported from `ggml/src/ggml-quants.c` @ `b10451` (`:891`, `:1869`, `:276`, over `make_qkx2_quants:799`, `make_qx_quants:628`, `nearest_int:621`) and gated **byte-identical** against a frozen golden captured from that reference compiled `-ffp-contract=off`. Two traps changed bytes and are recorded in the source rather than left to be re-found: `nearest_int` is the `+12582912.0` add-and-mask trick and rounds half to EVEN, not `round`; and C `roundf` in `quantize_row_q8_0_ref` rounds half AWAY FROM ZERO where `np.rint` rounds half to even, so `np.rint` mis-encodes every exact `.5` — caught by a crafted `[0.5, -0.5, 1.5, ...]` case, `[1,-1,2,-2,3,-3]` expected against `[0,0,2,-2,2,-2]` observed. **Exact arithmetic, from the converter's OWN type resolver run over the real topology** rather than bpw times a parameter count: 1719 output tensors carrying **313,890,512,702** parameters — 321.32B less the 7.43B MTP block, which independently confirms the skip is the 2.31% the spec measured. Q2_K (experts Q2_K, everything else Q6_K) **100.35 GiB** at 2.746 mixed bpw, Q6_K 239.89, Q8_0 310.67, BF16 584.67. Against 119.63 GiB at 128K context and one sequence: KV **1.43 GiB** (11 MLA layers x `kv_lora_rank` 512 x 2 B plus an 11 x 64 B indexer side cache = 11,968 B/token), KDA recurrent state **0.14 GiB** (64 x 128 x 128 x 4 B x 34) plus conv states, leaving **~17.7 GiB** — arithmetic, not measurement, and [#1963](https://github.com/mudler/vllm.cpp/issues/1963)/[#1966](https://github.com/mudler/vllm.cpp/issues/1966) mean W5 re-derives it from the runner. The Q2_K figure is **100.35 and not the spec's 102.6** because that table stated every arm including layer 45. The fallback ladder is `want -> Q8_0 -> F32` rather than `want -> F32`, because with an F32 fallback a FINER arm can come out LARGER than a coarser one and a size table must not have that property; on the real model the F32 sliver is 0.08 GiB (the depthwise conv kernels at `ne0=4`, the 4-D downsample and 5-D patch-embed). The tensor inventory is EXACT, not inferred: the real 8.4 MB `model.safetensors.index.json` (76,108 entries) and shard headers 2/32/62 were read by HTTP RANGE, payload never fetched, confirming the three separate `{q,k,v}_conv1d` convs, `hc_{attn,ffn}_{fn,base,scale}` flat on the layer with **no `hc_head.*` at any layer**, `indexer.k_norm.bias` present (settling LayerNorm-with-bias over RMSNorm), and `index_kpool_compress_{ape,gate}` on 12 layers. Refusals by name for every unimplemented arm: the i-quants because an imatrix needs a forward pass, a forward pass needs 181 GiB, and the dependency is **circular on this fleet** (O5); Q3_K/Q4_K/Q5_K because their encoders are unported and shipping an ungated encoder is worse than refusing (O8); `--keep-mtp` because nothing on the loader side reads an MTP tail (O2); a non-`glm5_next` config; and an FP8 tensor whose `weight_scale_inv` companion is missing, which would otherwise produce a loadable, wrong file. Gate: 50 assertions on a SYNTHETIC tiny-shape checkpoint with an independent in-test GGUF reader — no real checkpoint, no GPU, no C++ build. **NOT REACHED and disclosed:** `glm5next` has no `general.architecture` dispatch entry (`src/vllm/entrypoints/model_loader.cpp:1000`), so the file the converter writes is not loadable by this tree; W1 owns that wiring, tracked as **O9** under #1998. **Owed on close:** O7, artifact production — needs the 300-600 GiB checkpoint staged on local disk, explicit developer authority for the download, and a box with room for source and output at once, plus the sha256, recipe and peak RSS the spec's §Evidence requires | feature | +| [#2068](https://github.com/mudler/vllm.cpp/issues/2068) | `KV-DSV4-MULTICACHE` | **`GPUModelRunner` carries at most ONE attention group and ONE cache per layer, so the seven-group topology W2 publishes is refused instead of allocated, and `ModelForwardInput` has no third channel that could carry it.** W3 of [#1925](https://github.com/mudler/vllm.cpp/issues/1925), after W1 ([#1960](https://github.com/mudler/vllm.cpp/issues/1960)) landed `SlidingWindowMLASpec` and the four `MLAAttentionSpec` fields and W2 ([#1973](https://github.com/mudler/vllm.cpp/issues/1973)) made `MakeDeepseekV4KVCache` publish DeepSeek-V4-Flash's real **167 entries in seven groups**; spec [`specs/kv-dsv4-multicache.md`](specs/kv-dsv4-multicache.md) `### W3 design`. **Three parts, one defect.** (1) Group selection has two arms and no `else` (`src/vllm/v1/worker/gpu/runner.cpp:608-627`): `full_attn_group_id_` takes the FIRST non-eagle `kFullAttention`/`kMlaAttention` group, `gdn_group_id_` a `kMamba` group, and `kSlidingWindowMla`, `kSlidingWindow`, `kChunkedLocalAttention` and a SECOND `kMlaAttention` group match nothing; both ids are plain `int`s (`include/vllm/v1/worker/gpu/runner.h:591-592`). (2) The allocation loop allocates one buffer per HIDDEN LAYER, because `membership_by_name` is computed only inside `if (has_mamba_group)` (`runner.cpp:942-957`) and DeepSeek-V4 has no Mamba group, so it falls into `is_full_attn = !is_gdn`. MEASURED by W2 on a probe binary with the refusal disabled: two groups published, `full_attn_group_id = 0`, `gdn_group_id = -1`, `attn_kv().size() = 4` for 4 hidden layers, and the dropped `kSlidingWindowMla` group's own `page_size_bytes() = 37440` received nothing, with no exception and no message. The per-layer taxonomy that records the routing is a three-valued enum (`runner.h:379-383`) and cannot say `this layer has four caches`. (3) `ModelForwardInput` carries exactly two cache references (`include/vllm/model_executor/models/model_registry.h:344-345`); `attn_kv`'s positional convention cannot address a cache set keyed by (layer x role) the way upstream's `compilation_config.static_forward_context` is keyed by layer NAME. **Upstream has no equivalent generalization to port:** its runner allocates per registered NAME from the start (`vllm/v1/worker/gpu_model_runner.py:7785-7801` walks every `AttentionLayerBase`), so one C4A layer contributes four keys and there is no one-cache-per-layer assumption to relax; what is mirrored is the KEY, which is why the third channel carries the `static_forward_context` name beside each cache. Pin `5559679229bc961848b121ccdeaa8fa5d79bec98`. **FIXED** by entering a generalized path only when the published topology leaves groups over after the target attention group, the recurrent group and the `fa_draft` slot -- exactly the set W2's refusal already computes -- so every model shipping today runs the legacy path untouched, on the `per_layer_attn_specs` empty-means-unchanged contract (`include/vllm/v1/kv_cache_interface.h:538-556`). **Byte-neutrality re-derived by sweep at base `c714b0234` rather than inherited: 36 group-emplacement sites across 31 factory files -- 25 `FullAttentionSpec`, 7 `MLAAttentionSpec`, 3 `MambaSpec`, 1 `SlidingWindowMLASpec`**; by runtime shape 27 single-group factory FILES, `kimi_linear` (MLA+Mamba), `nemotron_h` (FA+Mamba), `qwen3_5_common` (FA+Mamba+`fa_draft` behind `if (num_spec > 0)`) and `deepseek_v4`. The count differs from W2's `34 across 32` because W2 itself added the two DeepSeek-V4 sites. **FILES, not architectures:** several single-group files back more than one `REGISTER_VLLM_MODEL` (`gemma4`/`gemma4_unified`, `olmo2`/`olmo3`, `llama_dense`/`internlm3_llama`, `muse_glimmer`/`muse_glimmer_mm`, the three parakeets, and `llama_model_embedding` reusing `MakeLlamaForCausalLMKVCache`), so counted by REGISTERED ARCHITECTURE it is **42 total = 34 single-group + 7 multi-group + 1 that publishes nothing** (`qwen4_exp`, whose spec factory throws by name). Only `deepseek_v4` has leftovers -- and **seven is its MAXIMUM, not its count**: `add_mla`/`add_swa_mla` return early on empty names, so the published count is a function of the checkpoint's `compress_ratios` and ranges over **1..7**; an all-`ratio == 1` config publishes ONE group, the SWA group, which is a `SlidingWindowMLASpec` that `full_attn_group_id_` never binds to and is therefore itself the leftover. **Every reachable DeepSeek-V4 config is multi-cache.** The full gate includes the SACRED `test_qwen35_paged_engine` regression. **Consequence stated rather than implied: a DeepSeek-V4 engine at `--block-size` 128 or 256 now CONSTRUCTS and allocates all 167 buffers, and its FIRST forward refuses**, because `ModelRegistry::Forward` will not let a multi-cache topology be silently ignored by a forward that discards `attn_kv`. That refusal is what keeps the third channel from landing dead: it reads the channel's payload -- cache count, group count, first layer name -- not its nullness. **W2's refusal is KEPT and still reaches four shapes:** a group whose layer names do not all resolve, a second `kMamba` group, a group whose spec is neither an `AttentionSpec` nor a `MambaSpec`, and an EAGLE group ([#2084](https://github.com/mudler/vllm.cpp/issues/2084), a review repair -- the first cut enumerated three and tested three, and an eagle group therefore passed the refusal and then got no buffer). **Not in scope, named so nobody takes them for this:** the uniform-`block_size` deferral in `HybridKVCacheCoordinator` (`src/vllm/v1/core/kv_cache_coordinator.cpp:340-346`, an `assert` and therefore inert under `NDEBUG`), the block-table geometry and the `KVBytesPerBlock` / `recurrent_state_bytes` budget helpers that still count one page per hidden layer, all W4; and the forward consuming the caches, W5. Both listed under `## Owed` | bug | | [#2076](https://github.com/mudler/vllm.cpp/issues/2076) | `KV-DSV4-MULTICACHE` | **`ENG-MOE-LOADSTREAM` cites `src/vllm/model_executor/models/model_registry.cpp:411` for `ModelSource::FromSafetensorsOwned`, which is 19 lines past the end of that file.** Found while gating [#2068](https://github.com/mudler/vllm.cpp/issues/2068) (`KV-DSV4-MULTICACHE` W3); PRE-EXISTING and not that wave's defect. The file has **392** lines at base `c714b0234` and the symbol is at **line 211** there; W3's own `#include` and blank line move it to **line 213**, which is the value the repaired citation carries. (198 is the distance from the symbol at W3's head to 411, not a distance to the end of the file.) `scripts/check-agent-record.py` sorts a bad citation into `broken` when the line is out of range and `stale` when the line exists and holds something else, and ratchets both against `scripts/record-anchor-baseline.json`, which stood at `stale: 28, broken: 6, total: 34` -- so an anchor inside its bucket's budget is invisible, and this one was part of that budget. **What made it visible is a general property of the ratchet rather than a coincidence:** #2068 adds 63 lines to `model_registry.cpp`, which pushes the file past line 411; the anchor did not change and its target did not move, but the line it names started EXISTING, so it moved from `broken` to `stale` and the checker fired twice at once -- `RECORD ANCHOR REGRESSION in bucket 'stale': 29 > baseline 28` AND `record-anchor baseline STALE in bucket 'broken': 5 < baseline 6`. **A broken anchor becomes a stale one as soon as anyone lengthens the file it points into**, and the change that lengthens the file is charged for it; a single total would have absorbed the move silently, which is the argument for the two buckets. FIXED IN FLOW: the citation is repaired to `:213` and the baseline lowered to `stale: 28, broken: 5, total: 33` in the same commit, which is what the checker's own message prescribes. Nothing else in that record moved -- `include/vllm/model_executor/models/model_registry.h:60` still holds the shards-owner field, and `src/vllm/entrypoints/model_loader.cpp:365` (`LoadFromDir`) stays in the `stale` budget as unrelated pre-existing drift, untouched rather than silently swept in | bug | | [#2084](https://github.com/mudler/vllm.cpp/issues/2084) | `KV-DSV4-MULTICACHE` | **An EAGLE `AttentionSpec` group on a multi-cache topology gets NO buffer and NO refusal.** `initialize_kv_cache` excludes an eagle group from `attn_group_ids_` (`src/vllm/v1/worker/gpu/runner.cpp`, `is_attention_spec && !group.is_eagle_group`), but W3's narrowed refusal loop never tested `is_eagle_group` -- so such a group kept `why` empty, passed the refusal, and then received no buffer because the allocation loop iterates `attn_group_ids_`. That is a SUBSET of the published topology allocated in silence, verbatim what the refusal's own message says it prevents; the code comment enumerated "three shapes" and this is a fourth. Demonstrated rather than reasoned: setting `is_eagle_group=true` on the indexer-key group in `MakeMultiCacheKvConfig` yields `REQUIRE( 9 == 10 )` -- nine caches allocated instead of ten, no message. **Not reachable today**: `is_eagle_group` is set true in exactly one place in the tree (`tests/vllm/v1/worker/test_runner.cpp`), on a config that is not a multi-cache topology, and no registered factory sets it; the multi-cache path also refuses at the forward until W5, so no wrong tokens were possible. Found while reviewing [#2068](https://github.com/mudler/vllm.cpp/issues/2068). FIXED IN FLOW with W3: the refusal gains an eagle clause and `test_runner`'s refusal case gains a subcase that is RED without it (`CHECK_THROWS_AS ... did NOT throw at all`, 5 assertions failed) and green with it. **Refusing is the direction rather than dropping the `!group.is_eagle_group` filter**, because allocating a draft group as an ordinary named cache decides how speculation shares a multi-cache topology, and that decision belongs to the wave that gates the speculative path | bug | | [#2085](https://github.com/mudler/vllm.cpp/issues/2085) | `KV-DSV4-MULTICACHE` | **The multi-cache `PagedKvCache` view geometry contradicts the page it is built over.** Each buffer is `num_blocks * spec->page_size_bytes()` while its `FaDims` view is built from `spec->block_size`; for a spec whose page derives from a `storage_block_size` the two disagree. DeepSeek-V4's C4A latent (`block_size` 256, `compress_ratio` 4, so `storage_block_size` 64) has a **37440**-byte page while the view declares `{num_blocks, 256, 512}` = **131072** bytes per block, 3.5x what the page holds. `CheckKvCacheShape` cannot see it: it compares the backend's declared shape against that same view metadata, so it measures self-consistency rather than agreement with the allocation. **INERT today** -- `ModelRegistry::Forward` refuses a multi-cache index before any kernel reads a view. Found while reviewing [#2068](https://github.com/mudler/vllm.cpp/issues/2068). NOT fixed in flow and OWED to **W5** with the store path, listed under `## Owed` in `.agents/specs/kv-dsv4-multicache.md`, because resolving it is entangled with two things W3 cannot settle: the `fp8_ds_mla` 584 B/token layout is not expressible in `PagedKvCache` at all, and `tests/vllm/v1/worker/test_runner.cpp` pins `block_size == 256` for that entry as a literal that the resolution may have to contradict. Given its own `## Owed` entry rather than folded into the W4 non-uniform-`block_size` item, which is about pool budgeting (`KVBytesPerBlock` counting one page per layer) where this is about the view a kernel would index off | bug | | [#2074](https://github.com/mudler/vllm.cpp/issues/2074) | `BACKEND-CUDA-SM087` | **`cmake/CudaArchFeatures.cmake`'s `fa2` row still labels its Ampere `sm_8x` cells "NO Ampere board ran them here", and an `sm_87` board ran the FA-2 path on 2026-08-26.** Jetson AGX Orin IS `sm_87`, one of the four cells that label covers. Row `MODEL-MM-dots3-note` ([#699](https://github.com/mudler/vllm.cpp/issues/699)) leased `orin:gpu0`, measured FA-2 ON three ways (`fa2: ENABLED for [87]`, `VLLM_CPP_FLASH_ATTN:BOOL=ON`, `VLLM_CPP_CUDA_FA2_COMPILED_ARCHS "87"`) and ran `test_ops_mla_prefill` to 2,931,678 assertions on the device against 329,772 with `CUDA_VISIBLE_DEVICES=""`, with the window biting at `gpu_win` vs `gpu_none` = 1.06055. Evidence `.agents/specs/dots3-note.md` §4.8; that record is CITED, not re-measured here. Scope: `8.7` only, so the label is now wrong for one of its four cells and right for `8.0`, `8.6` and `8.9`, which one line cannot carry. The identical claim also sits in `.agents/specs/cuda-arch-ampere-fastpath.md` WA-1 and the `BACKEND-CUDA-SM087` backend-matrix cell. Found by W4b-3a of #699 and filed rather than fixed, because re-labelling another row's verification state is that row's decision | record | -| [#2087](https://github.com/mudler/vllm.cpp/issues/2087) | `SPEC-DFLASH2` | **At every `c > 1` the DFlash2 draft leaves the paged CUDA-graph fast path and re-attends the WHOLE batch's context every decode step.** `GPUModelRunner::propose_drafts_block` is the only production caller of `ForwardBlockLogitsWithDeviceKV` and passes one store per proposing row, so `P == num proposing rows`; the fast path is gated on `P == 1` (`qwen3_dflash.cpp:1577`) and everything above it falls to `:1888-1930`, whose own comment says it is "not capture-targeted". That fallback materializes `2 x L` `[C, kdim]` context buffers, then `ForwardWithCtxKVDev` (`:664`) allocates `[Ncomb = C + Tq]` query and output buffers per layer (`:792-794`, `:811`) and calls `vt::DFlashBlockAttention` (`:818`), whose CUDA grid is over ALL `Ncomb` rows (`cuda_ops.cu:2634`, `:2643`, `:2650`) — an attention output computed for every context row of every request and then discarded at `:820-827`. Per step, per layer: `sum_r (ctx_r + 1 + k)^2` attention pairs instead of the paged route's `(1+k) x C`. It enters at c=2, grows with c, and is the shape of the measured stall (ours 60.25 -> 63.3 tok/s from c=4 to c=8 where vLLM goes 64.25 -> 80.0). Spec [`specs/dflash2-batch-propose.md`](../specs/dflash2-batch-propose.md) | perf | -| [#2088](https://github.com/mudler/vllm.cpp/issues/2088) | `SPEC-DFLASH2` | **A non-causal DFlash SWA layer runs with NO sliding window here, where upstream passes `per_layer_sliding_window` unconditionally.** Upstream resolves `(sliding_window, causal)` as two independent answers (`vllm/model_executor/models/qwen3_dflash.py:86-146` at pin `5559679229`) and hands the window to `Attention` with no reference to causality (`:229`), consuming `self.causal` one level out as attention metadata (`:234`, `:720`). Our tree conditions the window on causality in every kernel and in the paged seam's mask map — `qwen3_dflash_internal.h:125`, `cuda_ops.cu:1582`, `:1802`, `:1971`, `:2255`, `:2405`, `:2435`, `cpu_ops.cpp:2951`, `:3029` — so a declared `is_causal false` beside `sliding_attention` layers drops the window on all of them. That is every draft layer of the campaign subject, and it makes row 2 of the `.agents/specs/dflash2-draft-block-fa2.md` dispatch table unreachable from production. Acceptance-only, therefore invisible to a token gate, and the compute goes from `O(ctx x W)` to `O(ctx^2)` per query row. Owed: read the resolved `(causal, sliding_window)` pair off the draft's own `config.json` before writing code — the issue is INERT for a draft whose layers resolve `sliding_window == 0`. Listed under `## Owed` in [`specs/dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) | correctness | -| [#2089](https://github.com/mudler/vllm.cpp/issues/2089) | `SPEC-DFLASH2` | **The W11 draft-block route counters are blind to the production `P > 1` lane.** Both `NoteDflashBlockRoute` increments sit inside the `P == 1` branch (`qwen3_dflash.cpp:1487`, `:1502`); the materialized fallback at `:1888-1930` increments neither, so at every concurrency above one a route gate reads zero for both lanes while production runs a third route nothing names. #1890 put the counter inside the branch precisely so it would measure a capability rather than a class, and this is the hole that argument left. Listed under `## Owed` in [`specs/dflash2-spec-decode.md`](../specs/dflash2-spec-decode.md) | test-gap | +| [#2087](https://github.com/mudler/vllm.cpp/issues/2087) | `SPEC-DFLASH2` | **At every `c > 1` the DFlash2 draft leaves the paged CUDA-graph fast path and re-attends the WHOLE batch's context every decode step.** `GPUModelRunner::propose_drafts_block` is the only production caller of `ForwardBlockLogitsWithDeviceKV` and passes one store per proposing row, so `P == num proposing rows`; the fast path is gated on `P == 1` (`qwen3_dflash.cpp:1577`) and everything above it falls to `:1888-1930`, whose own comment says it is "not capture-targeted". That fallback materializes `2 x L` `[C, kdim]` context buffers, then `ForwardWithCtxKVDev` (`:664`) allocates `[Ncomb = C + Tq]` query and output buffers per layer (`:792-794`, `:811`) and calls `vt::DFlashBlockAttention` (`:818`), whose CUDA grid is over ALL `Ncomb` rows (`cuda_ops.cu:2634`, `:2643`, `:2650`) — an attention output computed for every context row of every request and then discarded at `:820-827`. Per step, per layer: `sum_r (ctx_r + 1 + k)^2` attention pairs instead of the paged route's `(1+k) x C`. It enters at c=2, grows with c, and is the shape of the measured stall (ours 60.25 -> 63.3 tok/s from c=4 to c=8 where vLLM goes 64.25 -> 80.0). Spec [`specs/dflash2-batch-propose.md`](specs/dflash2-batch-propose.md) | perf | +| [#2088](https://github.com/mudler/vllm.cpp/issues/2088) | `SPEC-DFLASH2` | **A non-causal DFlash SWA layer runs with NO sliding window here, where upstream passes `per_layer_sliding_window` unconditionally.** Upstream resolves `(sliding_window, causal)` as two independent answers (`vllm/model_executor/models/qwen3_dflash.py:86-146` at pin `5559679229`) and hands the window to `Attention` with no reference to causality (`:229`), consuming `self.causal` one level out as attention metadata (`:234`, `:720`). Our tree conditions the window on causality in every kernel and in the paged seam's mask map — `qwen3_dflash_internal.h:125`, `cuda_ops.cu:1582`, `:1802`, `:1971`, `:2255`, `:2405`, `:2435`, `cpu_ops.cpp:2951`, `:3029` — so a declared `is_causal false` beside `sliding_attention` layers drops the window on all of them. That is every draft layer of the campaign subject, and it makes row 2 of the `.agents/specs/dflash2-draft-block-fa2.md` dispatch table unreachable from production. Acceptance-only, therefore invisible to a token gate, and the compute goes from `O(ctx x W)` to `O(ctx^2)` per query row. Owed: read the resolved `(causal, sliding_window)` pair off the draft's own `config.json` before writing code — the issue is INERT for a draft whose layers resolve `sliding_window == 0`. Listed under `## Owed` in [`specs/dflash2-spec-decode.md`](specs/dflash2-spec-decode.md) | correctness | +| [#2089](https://github.com/mudler/vllm.cpp/issues/2089) | `SPEC-DFLASH2` | **The W11 draft-block route counters are blind to the production `P > 1` lane.** Both `NoteDflashBlockRoute` increments sit inside the `P == 1` branch (`qwen3_dflash.cpp:1487`, `:1502`); the materialized fallback at `:1888-1930` increments neither, so at every concurrency above one a route gate reads zero for both lanes while production runs a third route nothing names. #1890 put the counter inside the branch precisely so it would measure a capability rather than a class, and this is the hole that argument left. Listed under `## Owed` in [`specs/dflash2-spec-decode.md`](specs/dflash2-spec-decode.md) | test-gap | | [#2090](https://github.com/mudler/vllm.cpp/issues/2090) | `ENG-SCHED-CORE` | **The scheduler-mirror claim does not survive a spec-decode run.** `.agents/parity-ledger.md:467` scopes itself honestly to the waiting loop plus two defaults measured with no speculator, but `specs/scheduler-prefill-coschedule.md:132`, `specs/c8-p99-itl-tail-2026-07-18.md:28`, `:64-66` and `specs/async-serving.md:224` restate it unscoped. Two upstream mechanisms are absent and neither had an issue: `pad_spec_decode` (`vllm/v1/core/sched/scheduler.py:826-843`, `:1022-1025` at pin `5559679229`), which pads a newly admitted 1-token request to `1 + num_spec_tokens` to "preserve full cudagraph for this step", and the dynamic-SD lookup (`:1122-1125`, config-gated upstream). Both are recorded as deferrals in `include/vllm/v1/core/sched/scheduler.h:54-56`. It matters here because a ragged batch is a WHOLE-STEP cliff: `GraphEligibleQueryLen` (`src/vllm/v1/worker/gpu/cudagraph_dispatch.h:161-175`) refuses the entire step if any one request has `drafts + 1 != q`. INERT for the #1574 ladder, which runs `--no-enable-prefix-caching`; filed because it becomes live the moment prefix caching does. Verified in the same read: no `O(num_running^2)` term exists in our `schedule()` that upstream lacks | parity | -| [#2042](https://github.com/mudler/vllm.cpp/issues/2042) | `SPEC-DFLASH2` | **`--enable-prefix-caching` with a DFlash2 draft kills EngineCore on the first request that takes a cache hit, at concurrency 1, and that makes the SGLang-compat `lpm` scheduler unreachable.** Measured on `3d895a202`, `sm_121a`, dgx:gpu0 under an `rc` lease, DFlash2 k=8, 1024 in / 512 out: the same binary serves 8/8 with `--no-enable-prefix-caching` and reads `ok=0 failed=8` with it on, throwing `propose_drafts_block: context position discontinuity` from inside the EngineCore step, after which every later request returns `[request submitted to a stopped AsyncLLM]`. **The invariant is the DETECTOR, not the defect**, and the three facts that settle it are: the scheduler admits a cache-hit request with `num_computed_tokens` already equal to the cached prefix (`sched/scheduler.cpp`, the waiting-admission `get_computed_blocks` arm) and the worker turns that straight into absolute positions (`prepare_inputs.cpp`, `positions[t] = num_computed_tokens_cpu[r] + query_pos[t]`); the target is served from cache and never produces the aux hidden states the draft projects, so the private store genuinely holds ZERO context rows while the target has committed N, which the second `VT_CHECK` (`L == DeviceKVNumCtx`) confirms rather than contradicts; and **upstream never reaches that state because it keeps no private store at all** — its DFlash draft writes the context K/V into the engine's own paged KV cache through `attn.impl.do_kv_cache_update(...)` (`vllm/model_executor/models/qwen3_dflash.py:601-619` at pin `5559679229`) on a slot mapping built from the TARGET's block table (`vllm/v1/spec_decode/dflash.py:145-153`), so a prefix hit hands it the draft context for free. FIXED by mirroring upstream's OTHER answer, the one for a proposer that cannot serve a request: an EMPTY draft and the target running alone (`vllm/v1/spec_decode/ngram_proposer.py:150-159`, `suffix_decoding.py:55-62`, both `continue` and neither raises), which is [#1919](https://github.com/mudler/vllm.cpp/issues/1919)'s `disabled` fallback reached from a second place. **STACKED ON [#2010](https://github.com/mudler/vllm.cpp/pull/2010) ([#2008](https://github.com/mudler/vllm.cpp/issues/2008)) AND CANNOT LAND FIRST, for correctness rather than tidiness:** the classification keys on #2010's `first_sight` predicate ("this runner has never held context for this request"), and under the pre-#2010 row-indexed state that question could not be asked, because a request the batch had MOVED presented identically to a never-seen one — so the same fallback would have swallowed #2008's crash and turned it into a silent acceptance loss. Measured, not argued: mutation M3 drops the freshness gate and reddens exactly that assertion. #2010 does NOT fix this — on its head the engine still throws the discontinuity on the second request, seven times in one run. **What it buys and costs is stated rather than implied, and it is not a free win:** prefix caching's TTFT half is kept because the target still skips the cached prefill, and a hit request stops speculating for its life, so on a shared-system-prompt workload prefix caching and DFlash2 become mutually exclusive in effect and output throughput can fall; what is unambiguously fixed is that the configuration is currently a CRASH. The repair that removes the trade is the paged context store owed under [dflash2-ctx-store-capacity.md](../specs/dflash2-ctx-store-capacity.md) and tracked by #1919; a cheaper partial that keeps speculation over a TRUNCATED draft context anchored at the cache boundary is recorded under `## Owed` and deliberately not taken, because it moves draft acceptance and acceptance cannot be measured without a device. Gated by `tests/vllm/v1/spec_decode/test_dflash2_prefix_cache.cpp` (5 cases, 69 assertions, CPU, through the production `AsyncLLM` front): red-before 3/5 cases fail with the engine dead, green-after 5/5, with G3 and G5 green on both sides as controls. Wave spec [dflash2-prefix-cache.md](../specs/dflash2-prefix-cache.md) | bug | -| [#2067](https://github.com/mudler/vllm.cpp/issues/2067) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The `glm5_next` converter writes a file nothing in this tree can open: register the architecture, and give it its `general.architecture` dispatch row (O9).** W1 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998); spec [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) §W1. W7a ([#2011](https://github.com/mudler/vllm.cpp/issues/2011)) authored `scripts/convert-glm5-next-gguf.py`, which emits `general.architecture = glm5next`; that key had no row in `kGgufArchArms` and `Glm5NextForConditionalGeneration` was registered by no translation unit, so both entry points refused the model by name as unrecognized and every downstream wave (W3, W5, W6, W7b) had nothing to load. **One parser, two sources.** `Glm5NextHfConfigFromGguf` reads the converter's metadata and synthesizes an HF-shaped `text_config`/`vision_config` under the *same key spellings* `config.json` uses, so a GGUF descends through the SAME `ParseGlm5NextParams` a `config.json` does — one validation surface, not two that can drift. `ParseGlm5NextParams` mirrors `Glm5NextTextConfig.__post_init__` and all five `validate_architecture` rejections at transformers **v5.16.1** (`eb4d9e2a64`, the first release carrying `glm5_next`; `v5.16.0` is 404): the `full_attention` -> `deepseek_sparse_attention` layer-kind rewrite (so `Glm5NextLayerKind` has no `kFullAttention` enumerator at all and the checkpoint's spelling is unrepresentable rather than merely unused); the `linear_attn_config` -> `linear_{head_dim,num_heads,conv_kernel_dim,lower_bound}` remap together with its `safe_gate`-defaults-True rule, and the deliberate IGNORING of that dict's `kda_layers`/`full_attn_layers` index lists, which the reference never reads; the `mlp_layer_types` default `[dense]*min(3,L) + [sparse]*(L-3)`; the `indexer_types` freq/offset schedule; and the forced `head_dim = qk_rope_head_dim`, `qk_head_dim = qk_rope_head_dim + qk_nope_head_dim` overrides. **The two validators are exact complements, and that is the structural finding.** Upstream RAISES when `qk_rope_head_dim > 0` ("Expecting NoPE for the DSA attention layers"); our `MlaBlockDims::Validate` RAISES when it is not `> 0` (`mla_attention.cpp:90-93`). No value satisfies both. W1 mirrors upstream and accepts `0`; the relaxation is W3's and is recorded as **O11**, pinned by a test so W3 cannot land the geometry without moving the pin. The HF->GGUF tensor name map is enumerated structurally per layer KIND, and the config builder uses it for one reachable, shard-safe check: a `blk.N` that carries KDA tensors while the metadata declares that layer `deepseek_sparse_attention` (or the converse) is refused, because absence proves nothing on a sharded file but a CONTRADICTION is a wrong model loading quietly. **Scope honesty.** This makes the architecture RESOLVE and its config PARSE and VALIDATE. It does NOT make the model load and it does NOT make it forward: the loader, the forward and the KV-cache spec each refuse by name, naming the missing primitive and the wave that owes it (**O10**). No token, no speed, no artifact — O1 holds unchanged, and no oracle can execute this model on any device this fleet reaches. vLLM implements `glm5_next` at NO revision, so no pin was advanced and none is owed; the sole admissible reference is transformers, and **W0's lane pin for `v5.16.1` is still unwritten** — this wave cites the revision it read without recording a pin, which stays W0's deliverable | feature | +| [#2042](https://github.com/mudler/vllm.cpp/issues/2042) | `SPEC-DFLASH2` | **`--enable-prefix-caching` with a DFlash2 draft kills EngineCore on the first request that takes a cache hit, at concurrency 1, and that makes the SGLang-compat `lpm` scheduler unreachable.** Measured on `3d895a202`, `sm_121a`, dgx:gpu0 under an `rc` lease, DFlash2 k=8, 1024 in / 512 out: the same binary serves 8/8 with `--no-enable-prefix-caching` and reads `ok=0 failed=8` with it on, throwing `propose_drafts_block: context position discontinuity` from inside the EngineCore step, after which every later request returns `[request submitted to a stopped AsyncLLM]`. **The invariant is the DETECTOR, not the defect**, and the three facts that settle it are: the scheduler admits a cache-hit request with `num_computed_tokens` already equal to the cached prefix (`sched/scheduler.cpp`, the waiting-admission `get_computed_blocks` arm) and the worker turns that straight into absolute positions (`prepare_inputs.cpp`, `positions[t] = num_computed_tokens_cpu[r] + query_pos[t]`); the target is served from cache and never produces the aux hidden states the draft projects, so the private store genuinely holds ZERO context rows while the target has committed N, which the second `VT_CHECK` (`L == DeviceKVNumCtx`) confirms rather than contradicts; and **upstream never reaches that state because it keeps no private store at all** — its DFlash draft writes the context K/V into the engine's own paged KV cache through `attn.impl.do_kv_cache_update(...)` (`vllm/model_executor/models/qwen3_dflash.py:601-619` at pin `5559679229`) on a slot mapping built from the TARGET's block table (`vllm/v1/spec_decode/dflash.py:145-153`), so a prefix hit hands it the draft context for free. FIXED by mirroring upstream's OTHER answer, the one for a proposer that cannot serve a request: an EMPTY draft and the target running alone (`vllm/v1/spec_decode/ngram_proposer.py:150-159`, `suffix_decoding.py:55-62`, both `continue` and neither raises), which is [#1919](https://github.com/mudler/vllm.cpp/issues/1919)'s `disabled` fallback reached from a second place. **STACKED ON [#2010](https://github.com/mudler/vllm.cpp/pull/2010) ([#2008](https://github.com/mudler/vllm.cpp/issues/2008)) AND CANNOT LAND FIRST, for correctness rather than tidiness:** the classification keys on #2010's `first_sight` predicate ("this runner has never held context for this request"), and under the pre-#2010 row-indexed state that question could not be asked, because a request the batch had MOVED presented identically to a never-seen one — so the same fallback would have swallowed #2008's crash and turned it into a silent acceptance loss. Measured, not argued: mutation M3 drops the freshness gate and reddens exactly that assertion. #2010 does NOT fix this — on its head the engine still throws the discontinuity on the second request, seven times in one run. **What it buys and costs is stated rather than implied, and it is not a free win:** prefix caching's TTFT half is kept because the target still skips the cached prefill, and a hit request stops speculating for its life, so on a shared-system-prompt workload prefix caching and DFlash2 become mutually exclusive in effect and output throughput can fall; what is unambiguously fixed is that the configuration is currently a CRASH. The repair that removes the trade is the paged context store owed under [dflash2-ctx-store-capacity.md](specs/dflash2-ctx-store-capacity.md) and tracked by #1919; a cheaper partial that keeps speculation over a TRUNCATED draft context anchored at the cache boundary is recorded under `## Owed` and deliberately not taken, because it moves draft acceptance and acceptance cannot be measured without a device. Gated by `tests/vllm/v1/spec_decode/test_dflash2_prefix_cache.cpp` (5 cases, 69 assertions, CPU, through the production `AsyncLLM` front): red-before 3/5 cases fail with the engine dead, green-after 5/5, with G3 and G5 green on both sides as controls. Wave spec [dflash2-prefix-cache.md](specs/dflash2-prefix-cache.md) | bug | +| [#2067](https://github.com/mudler/vllm.cpp/issues/2067) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The `glm5_next` converter writes a file nothing in this tree can open: register the architecture, and give it its `general.architecture` dispatch row (O9).** W1 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998); spec [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) §W1. W7a ([#2011](https://github.com/mudler/vllm.cpp/issues/2011)) authored `scripts/convert-glm5-next-gguf.py`, which emits `general.architecture = glm5next`; that key had no row in `kGgufArchArms` and `Glm5NextForConditionalGeneration` was registered by no translation unit, so both entry points refused the model by name as unrecognized and every downstream wave (W3, W5, W6, W7b) had nothing to load. **One parser, two sources.** `Glm5NextHfConfigFromGguf` reads the converter's metadata and synthesizes an HF-shaped `text_config`/`vision_config` under the *same key spellings* `config.json` uses, so a GGUF descends through the SAME `ParseGlm5NextParams` a `config.json` does — one validation surface, not two that can drift. `ParseGlm5NextParams` mirrors `Glm5NextTextConfig.__post_init__` and all five `validate_architecture` rejections at transformers **v5.16.1** (`eb4d9e2a64`, the first release carrying `glm5_next`; `v5.16.0` is 404): the `full_attention` -> `deepseek_sparse_attention` layer-kind rewrite (so `Glm5NextLayerKind` has no `kFullAttention` enumerator at all and the checkpoint's spelling is unrepresentable rather than merely unused); the `linear_attn_config` -> `linear_{head_dim,num_heads,conv_kernel_dim,lower_bound}` remap together with its `safe_gate`-defaults-True rule, and the deliberate IGNORING of that dict's `kda_layers`/`full_attn_layers` index lists, which the reference never reads; the `mlp_layer_types` default `[dense]*min(3,L) + [sparse]*(L-3)`; the `indexer_types` freq/offset schedule; and the forced `head_dim = qk_rope_head_dim`, `qk_head_dim = qk_rope_head_dim + qk_nope_head_dim` overrides. **The two validators are exact complements, and that is the structural finding.** Upstream RAISES when `qk_rope_head_dim > 0` ("Expecting NoPE for the DSA attention layers"); our `MlaBlockDims::Validate` RAISES when it is not `> 0` (`mla_attention.cpp:90-93`). No value satisfies both. W1 mirrors upstream and accepts `0`; the relaxation is W3's and is recorded as **O11**, pinned by a test so W3 cannot land the geometry without moving the pin. The HF->GGUF tensor name map is enumerated structurally per layer KIND, and the config builder uses it for one reachable, shard-safe check: a `blk.N` that carries KDA tensors while the metadata declares that layer `deepseek_sparse_attention` (or the converse) is refused, because absence proves nothing on a sharded file but a CONTRADICTION is a wrong model loading quietly. **Scope honesty.** This makes the architecture RESOLVE and its config PARSE and VALIDATE. It does NOT make the model load and it does NOT make it forward: the loader, the forward and the KV-cache spec each refuse by name, naming the missing primitive and the wave that owes it (**O10**). No token, no speed, no artifact — O1 holds unchanged, and no oracle can execute this model on any device this fleet reaches. vLLM implements `glm5_next` at NO revision, so no pin was advanced and none is owed; the sole admissible reference is transformers, and **W0's lane pin for `v5.16.1` is still unwritten** — this wave cites the revision it read without recording a pin, which stays W0's deliverable | feature | | [#2070](https://github.com/mudler/vllm.cpp/issues/2070) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The shared config reader synthesizes `layer_types` from `linear_attn_config.kda_layers` as ONE-INDEXED, and GLM-5.3-Flash's list is ZERO-INDEXED.** Found while implementing W1 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998) ([#2067](https://github.com/mudler/vllm.cpp/issues/2067)) by a test that erased `layer_types` to check the port reproduces upstream's default schedule; it produced a schedule off by one and every assertion about which layer is KDA failed. `src/vllm/transformers_utils/hf_config.cpp` synthesizes `cfg.layer_types` from `text_config.linear_attn_config.kda_layers` when `layer_types` is absent, resolving it as `is_kda[one_indexed - 1] = true` and dropping any entry below 1 — correct for Kimi-Linear, whose upstream defines `is_kda_layer(l) := (l+1) in kda_layers`. `zai-org/GLM-5.3-Flash`'s list is ZERO-indexed, and the checkpoint settles it two ways: it contains `0`, which a one-indexed list of 45 layers cannot, and its maximum is `44` on `num_hidden_layers: 45`. Read through the one-indexed rule the `0` is dropped and everything else shifts down, so layer 2 comes out `full_attention` where the checkpoint calls it `linear_attention` — a wrong attention kind on a third of the stack, chosen silently. **Worse than an ordinary off-by-one:** the transformers reference IGNORES `kda_layers` entirely for `glm5_next`. `Glm5NextTextConfig.__post_init__` reads only `head_dim`, `num_heads`, `short_conv_kernel_size` and `gate_lower_bound` out of that dict and derives the schedule from the top-level `layer_types` or from `idx % 4 != 3`, so the shared reader would be deriving a load-bearing schedule from a list upstream never consults, under another family's indexing convention. **Not live on `main` today**, and that is the only reason this is not a shipped defect: no `glm5_next` reached `ParseHfConfig` at all until #2067 registered it, and every published `glm5_next` config carries an explicit `layer_types`, which the synthesis is guarded behind (`cfg.layer_types.empty()`). It is a trap set for the first wave to hand this model a config without one — which is what a converter, a hand-written test fixture, or a text-only variant produces. **REPAIRED IN FLOW by #2067:** `ParseGlm5NextParams` resolves `layer_types` from its own `text_config` and from upstream's `idx % 4 != 3` default, never from `cfg.layer_types`, so this model's schedule cannot be decided by a heuristic written for another family; it additionally cross-checks `kda_layers` / `full_attn_layers` against the resolved schedule AS ZERO-INDEXED and refuses on a disagreement rather than picking a winner. The shared reader's Kimi-Linear branch is left exactly as it is — it is correct for the family it was written for, and narrowing it is a change to Kimi-Linear's behaviour this row has no gate for. Pinned by `test_glm5_next_scaffold.cpp`'s `kda_layers is ZERO-indexed, and the schedule ignores it` case; mutation M5, taking `layer_types` from the shared reader again, reds it | bug | -| [#2097](https://github.com/mudler/vllm.cpp/issues/2097) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash's KDA forget gate takes the OTHER branch, and reusing `kimi_kda.cpp`'s would have produced a fluent wrong model on 34 of 45 layers.** W2 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998); spec [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) §W2; claim `CLAIM-GLM53-FLASH-W2`. `Glm5NextTextForgetGate.forward` branches on `safe_gate_lower_bound`, and `zai-org/GLM-5.3-Flash` sets `linear_attn_config.gate_lower_bound: -5.0`, so it computes `-5.0 * sigmoid(exp(A_log) * (f_b(f_a(x)) + dt_bias))` (`modular_glm5_next.py:398-399` @ transformers **v5.16.1**, the lane pin W0 ([#2096](https://github.com/mudler/vllm.cpp/issues/2096)) owes; anchors re-read at that tag and the file's 95,314 bytes confirmed). Our Kimi-Linear KDA (`src/vllm/model_executor/models/kimi_kda.cpp:60`) implements the ELSE branch, `-exp(A_log) * softplus(g + dt_bias)`. **Three things differ and each one alone keeps the model fluent:** the SHAPE (bounded logistic against unbounded softplus), the RANGE (the sigmoid branch cannot leave `[bound, 0]`, so the per-step decay is floored at `exp(-5)`), and the SIGN of `decay_rate` — `+exp(A_log)` multiplies `g` INSIDE the sigmoid where the softplus branch negates it OUTSIDE, so reusing one sign for the other mirrors the gate about `g = 0` and turns "forgets nothing" into "forgets everything" with no NaN anywhere. **RED-FIRST, captured before any implementation existed:** the sigmoid branch's hand-derived values asserted against `KdaDecayGate` gave `-0.693147` for `-2.5`, `-100` for `-5` and `-3.13262` for `-0.000227` — 3 cases, 8 assertions, 5 failed. That discriminator is now permanent rather than a one-off capture: on a shared randomized draw the two branches disagree at relative L2 **0.863** while our own softplus fallback agrees with `kimi_kda`'s below 1e-6, which separates "different branch" from "broken port", and Kimi-Linear's answer leaves `[-5, 0]` on that very draw so no rescaling of one produces the other. Landed as `glm5_next_kda.{h,cpp}`: the forget gate (BOTH branches, because `linear_lower_bound` is optional upstream, typed `float` or `None`), the low-rank bottleneck, the strict-fp32 `Glm5NextRmsNormGated` (`:414-426`), `l2norm` (`:429-437`), the conv-weight concatenation and the assembled host layer (`:641-746`). **Three traps gated as their own cases, each a plausible default that is not this model's value:** the output norm's activation is SIGMOID and not silu (`:412`); its eps is `rms_norm_eps` 1e-5 PASSED IN at `:635`, not the 1e-6 constructor default, which is a 2.3x error on the near-zero-variance row that is the only place an eps does anything; and `l2norm`'s eps is ADDED INSIDE the root, not `F.normalize`'s `max(norm, eps)`, which agrees to 1e-7 on an ordinary row and returns 0.0995 against 1.0 on a near-zero one — the row a randomized test never draws. **The three layout facts are gated, not commented:** the checkpoint's three separate `{q,k,v}_conv1d` depthwise convs concatenate in **q, k, v** order, gated against references built with the WRONG pairing — a `[q; k; v]` stream convolved with `[q_w; v_w; k_w]` and with `[k_w; q_w; v_w]` — because merely swapping two weight tensors moves the answer under ANY fixed order and so gates nothing; `g`, `beta` and the output gate read the **PRE-conv** hidden states (`:709`, `:710`, `:742`), proven visible by a reference variant that fuses them; and the cache is a `[conv_dim, K]` conv state plus an **fp32** `[H, Dk, Dv]` recurrent state, gated by a 6-then-1 split reproducing the 7-token one-shot forward while a fresh cache does not. `dt_bias` is REQUIRED and refused by name when absent or misshaped: `:384` declares it unconditionally and `:393` always adds it, so an optional mode would silently compute a different gate that stays finite and plausible. **Two annotated `f32` exceptions, each with upstream's own reason:** the norm is strict fp32 and does not downcast the weights (`:417-421`), and the recurrent state is cast to fp32 explicitly at `:739` because `:452` says the states are the rounding-sensitive part. **O14 — `vt::KdaChunkPrefill` cannot serve this model:** it FUSES `-exp(a_log)*softplus(g_raw+dt_bias)` in the vendored FLA cubins and in its CPU reference (`src/vt/cpu/cpu_ops.cpp:1779-1786`), and inverting that needs `g_raw = log(exp(-target)-1)`, which diverges as the gate approaches 0 — where most channels of 34 layers sit. Both prefill and decode therefore route through `vt::KdaGatedDeltaRule`, which is branch-agnostic because it consumes an already-computed log-decay. **O15 — NOT REACHED and disclosed:** `Glm5NextForConditionalGeneration::Forward` still refuses by name (O10), so at this merge commit the only call sites are the focused gate's; W5, the assembled text forward, owns the wiring on this row and has no issue of its own, so #1998 tracks it. Gate: CPU only, 28 cases / **342 assertions**, no GPU, no checkpoint, no pin advance, and no claim about the MODEL — O1 holds unchanged and nothing on this fleet can execute the reference | feature | -| [#2096](https://github.com/mudler/vllm.cpp/issues/2096) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash gates against `transformers` and its lane-scoped `v5.16.1` pin is unwritten: write it.** No oracle other than `transformers` implements `glm5_next` -- vLLM implements nothing at the parity pin `555967922` or at `main`, and vllm#53906 is OPEN and therefore inadmissible. `transformers` carries the architecture from `eb4d9e2a64a0` (transformers#48342, merged 2026-08-26T14:26:41Z) and the FIRST release carrying it is `v5.16.1`, bounded rather than assumed: `modeling_glm5_next.py` is HTTP 200 at `v5.16.1`, 404 at `v5.16.0` and 404 at `v5.15.1`, re-measured 2026-08-27. The registry pin is `5.14.1` and does not contain `Glm5Next`, so this row needs a lane-scoped second pin with `gateable = no`, expiring when vLLM registers `glm5_next`. Discharges O12 in [glm5-next-flash.md](../specs/glm5-next-flash.md). W0 of campaign issue #1998 | record | +| [#2097](https://github.com/mudler/vllm.cpp/issues/2097) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash's KDA forget gate takes the OTHER branch, and reusing `kimi_kda.cpp`'s would have produced a fluent wrong model on 34 of 45 layers.** W2 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998); spec [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) §W2; claim `CLAIM-GLM53-FLASH-W2`. `Glm5NextTextForgetGate.forward` branches on `safe_gate_lower_bound`, and `zai-org/GLM-5.3-Flash` sets `linear_attn_config.gate_lower_bound: -5.0`, so it computes `-5.0 * sigmoid(exp(A_log) * (f_b(f_a(x)) + dt_bias))` (`modular_glm5_next.py:398-399` @ transformers **v5.16.1**, the lane pin W0 ([#2096](https://github.com/mudler/vllm.cpp/issues/2096)) owes; anchors re-read at that tag and the file's 95,314 bytes confirmed). Our Kimi-Linear KDA (`src/vllm/model_executor/models/kimi_kda.cpp:60`) implements the ELSE branch, `-exp(A_log) * softplus(g + dt_bias)`. **Three things differ and each one alone keeps the model fluent:** the SHAPE (bounded logistic against unbounded softplus), the RANGE (the sigmoid branch cannot leave `[bound, 0]`, so the per-step decay is floored at `exp(-5)`), and the SIGN of `decay_rate` — `+exp(A_log)` multiplies `g` INSIDE the sigmoid where the softplus branch negates it OUTSIDE, so reusing one sign for the other mirrors the gate about `g = 0` and turns "forgets nothing" into "forgets everything" with no NaN anywhere. **RED-FIRST, captured before any implementation existed:** the sigmoid branch's hand-derived values asserted against `KdaDecayGate` gave `-0.693147` for `-2.5`, `-100` for `-5` and `-3.13262` for `-0.000227` — 3 cases, 8 assertions, 5 failed. That discriminator is now permanent rather than a one-off capture: on a shared randomized draw the two branches disagree at relative L2 **0.863** while our own softplus fallback agrees with `kimi_kda`'s below 1e-6, which separates "different branch" from "broken port", and Kimi-Linear's answer leaves `[-5, 0]` on that very draw so no rescaling of one produces the other. Landed as `glm5_next_kda.{h,cpp}`: the forget gate (BOTH branches, because `linear_lower_bound` is optional upstream, typed `float` or `None`), the low-rank bottleneck, the strict-fp32 `Glm5NextRmsNormGated` (`:414-426`), `l2norm` (`:429-437`), the conv-weight concatenation and the assembled host layer (`:641-746`). **Three traps gated as their own cases, each a plausible default that is not this model's value:** the output norm's activation is SIGMOID and not silu (`:412`); its eps is `rms_norm_eps` 1e-5 PASSED IN at `:635`, not the 1e-6 constructor default, which is a 2.3x error on the near-zero-variance row that is the only place an eps does anything; and `l2norm`'s eps is ADDED INSIDE the root, not `F.normalize`'s `max(norm, eps)`, which agrees to 1e-7 on an ordinary row and returns 0.0995 against 1.0 on a near-zero one — the row a randomized test never draws. **The three layout facts are gated, not commented:** the checkpoint's three separate `{q,k,v}_conv1d` depthwise convs concatenate in **q, k, v** order, gated against references built with the WRONG pairing — a `[q; k; v]` stream convolved with `[q_w; v_w; k_w]` and with `[k_w; q_w; v_w]` — because merely swapping two weight tensors moves the answer under ANY fixed order and so gates nothing; `g`, `beta` and the output gate read the **PRE-conv** hidden states (`:709`, `:710`, `:742`), proven visible by a reference variant that fuses them; and the cache is a `[conv_dim, K]` conv state plus an **fp32** `[H, Dk, Dv]` recurrent state, gated by a 6-then-1 split reproducing the 7-token one-shot forward while a fresh cache does not. `dt_bias` is REQUIRED and refused by name when absent or misshaped: `:384` declares it unconditionally and `:393` always adds it, so an optional mode would silently compute a different gate that stays finite and plausible. **Two annotated `f32` exceptions, each with upstream's own reason:** the norm is strict fp32 and does not downcast the weights (`:417-421`), and the recurrent state is cast to fp32 explicitly at `:739` because `:452` says the states are the rounding-sensitive part. **O14 — `vt::KdaChunkPrefill` cannot serve this model:** it FUSES `-exp(a_log)*softplus(g_raw+dt_bias)` in the vendored FLA cubins and in its CPU reference (`src/vt/cpu/cpu_ops.cpp:1779-1786`), and inverting that needs `g_raw = log(exp(-target)-1)`, which diverges as the gate approaches 0 — where most channels of 34 layers sit. Both prefill and decode therefore route through `vt::KdaGatedDeltaRule`, which is branch-agnostic because it consumes an already-computed log-decay. **O15 — NOT REACHED and disclosed:** `Glm5NextForConditionalGeneration::Forward` still refuses by name (O10), so at this merge commit the only call sites are the focused gate's; W5, the assembled text forward, owns the wiring on this row and has no issue of its own, so #1998 tracks it. Gate: CPU only, 28 cases / **342 assertions**, no GPU, no checkpoint, no pin advance, and no claim about the MODEL — O1 holds unchanged and nothing on this fleet can execute the reference | feature | +| [#2096](https://github.com/mudler/vllm.cpp/issues/2096) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash gates against `transformers` and its lane-scoped `v5.16.1` pin is unwritten: write it.** No oracle other than `transformers` implements `glm5_next` -- vLLM implements nothing at the parity pin `555967922` or at `main`, and vllm#53906 is OPEN and therefore inadmissible. `transformers` carries the architecture from `eb4d9e2a64a0` (transformers#48342, merged 2026-08-26T14:26:41Z) and the FIRST release carrying it is `v5.16.1`, bounded rather than assumed: `modeling_glm5_next.py` is HTTP 200 at `v5.16.1`, 404 at `v5.16.0` and 404 at `v5.15.1`, re-measured 2026-08-27. The registry pin is `5.14.1` and does not contain `Glm5Next`, so this row needs a lane-scoped second pin with `gateable = no`, expiring when vLLM registers `glm5_next`. Discharges O12 in [glm5-next-flash.md](specs/glm5-next-flash.md). W0 of campaign issue #1998 | record | | [#2099](https://github.com/mudler/vllm.cpp/issues/2099) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **`scripts/check-oracle-pins.py` never parses an `oracle-pin-lane` block, so both lane pins in `.agents/oracles/transformers.md` are unchecked prose.** Its `BLOCK` regex is `^```oracle-pin\n`, and the newline means an `oracle-pin-lane` fence never matches; a repository-wide search for `oracle-pin-lane` returns the block itself and one prose reference, so nothing reads it. MEASURED on `row/MODEL-MM-GLM53-FLASH-W0`: corrupting the `glm5_next` lane `pin`, `gateable` or `pinned_on`, and deleting the lane block outright, each leave the checker at exit 0, while the same corruption of the registry `oracle-pin` block reds it. Every rule the registry gate holds is therefore unenforced on a lane pin, and the checker's `--self-test` corpus and `tests/scripts/test_check_oracle_pins.py` name no lane case. Found by W0 (#2096) while verifying its own gate; recorded as O13 rather than repaired, because W0's scope excludes every checker and the fix is a semantic checker change that owes a spec, a red-before mutation, and a decision about which keys a lane record requires | gap | | [#2101](https://github.com/mudler/vllm.cpp/issues/2101) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **`main` does not compile under MSVC: seven range-`for` loop variables named `n` in `Glm5NextExpectedGgufTensors` hide the function-scope `const size_t n`, and `/W4 /WX` turns C4456 into `error C2220`.** Landed by [#2067](https://github.com/mudler/vllm.cpp/issues/2067) (PR [#2073](https://github.com/mudler/vllm.cpp/pull/2073), commit `47a2b35a5`), which authored `src/vllm/model_executor/models/glm5_next_weights.cpp` whole. `windows-msvc-cpu` and `windows-msvc-vulkan` both fail the build, so **every pull request that merges current `main` inherits a red Windows pair** — and it looks exactly like the long-standing [#584](https://github.com/mudler/vllm.cpp/issues/584) crash in `gh pr checks`, same two job names and same red, distinguishable only by reading the log: #584 carries exit `-1073740791` and zero `error C####`, this carries one `error C2220` and no crash code. **The issue's stated cause is not the mechanism, and the correction matters for the fix.** Sibling scopes do not hide one another, so the five loops named in #2101 do not shadow each other; every one of them shadows `const size_t n` at `glm5_next_weights.cpp:252`, the layer-count local the function's own bounds check uses. Renaming loop variables to be distinct *from each other* would therefore have left the defect in place. **CI reported four sites and there are seven**, because MSVC stops at the first `error C2220`: 276, 279, 285 and 287 reached the log; 288, 293 and 299 never did. Found with GCC's `-Wshadow`, whose `shadows a previous local` diagnostic is the exact analogue of C4456 and which names all seven at once — red-before `rc=1`, green-after `rc=0` on the same command. Fixed by naming the function-scope local `layer_count` for what it is, which removes all seven hidings at their source, and by naming the seven loop variables `tn` so no bland one-letter name can collide there again. No pragma, no `/WX` relaxation, no suppression: the warning is correct. Green-after for the Windows build itself is the CI job, which cannot be run on this fleet | bug | | [#2003](https://github.com/mudler/vllm.cpp/issues/2003) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **The #1604 flip premise inverted at post-W2c `21fe11cf1`: `VT_TT_HOST_FREE_DECODE=0` (host-hybrid) outperforms the shipped eager DEFAULT 1.24x on the P150** — Qwen3-0.6B b1 greedy, order-alternated pairs ×3, in-process run 1 discarded, one `$HOME/gpu.lock` hold, `tt-smi -r` first: default median 10.822 tok/s (n=12, 10.51–11.03) vs opt-out median 13.369 (n=12; ≥13.09 on 11 of 12). The default arm is UNCHANGED against its 2026-08-21 figures (10.94–11.06 at `b86e3705f`), so what moved is a ~2.5x improvement of the opt-out arm whose mechanism is unattributed; the next traceable step is a per-op delta of the host-hybrid path from `b86e3705f` to `21fe11cf1`. Stated rather than implied: NO clock window was sampled (`tools/bench/gpu_clock_state.py` is NVIDIA-only), so every figure including the record entry that cites this issue is clock-unattributed and quotable only as such; one model shape, one board (Blackhole P150, aarch64 host, tt-metal `a3d33028975`); the captured opt-in arm was NOT retested (#1625 still blocks multi-request capture, #1627 still open). The shipped default now serves the slower of the two eager arms, which any gate using the default as denominator inherits | perf | | [#2110](https://github.com/mudler/vllm.cpp/issues/2110) | `MODEL-MM-QWEN4-EXP` | **W5b-1: `GdnBlockPaged` was sealed in `qwen3_5.cpp`'s anonymous namespace, so the `qwen4_exp` forward could not reach the block 36 of its 48 layers ARE.** `Qwen4ExpLayerKind` has only `kLinearAttention` and `kQwenSparseAttention`, so of the five blocks in that anonymous namespace the forward needs exactly one, and the MoE half was already reachable through `RunMoeBlock` (`include/vllm/model_executor/models/qwen3_5_moe_block.h`, landed by `ENG-HYBRID-PLACEMENT` in `f730eb11c`). Exposed by MIRRORING that precedent rather than relocating the block: `include/vllm/model_executor/models/qwen3_5_gdn_block.h` declares `RunGdnBlockPaged` + `BuildGdnStepInputs` over primitive `vt::` types, and the two definitions sit beside `RunMoeBlock` in `qwen3_5.cpp` outside the anonymous namespace. The `dense_attn_block.h` relocation shape was REJECTED for this block, and measured rather than asserted: `GdnBlockPaged` is 475 lines and delegates to `GdnBlockPagedMixedSpec`'s further 233, and a scan of the anonymous namespace's top-level definitions finds at least 28 sibling symbols those two read (`ProjectGdnQkvz`, `ProjectGdnBA`, `MatmulFp8CutlassD`, `GatherStateF32`, `StepDevInputs`, the eight env-flag readers, ...), so a verbatim header move drags a large part of an 11.6k-line file two other agents are concurrently editing, while the wrapper adds 39 lines and moves nothing. The per-step device upload (`StepDevInputs`) stays a SEPARATE opaque handle on purpose: `qwen3_5.cpp` builds it once per step and shares it across every layer, and a wrapper that rebuilt it inside each block call would reinstate the per-layer upload that was removed, 36 times per step on this architecture. Gated by a new case in `tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp` (`GDN block seam (CPU)`, +22 assertions at both gate dims) that drives the seam from primitive `vt::` types in a foreign TU and compares output AND mutated SSM/conv state bit-for-bit against `GdnBlockPagedForTest`. Mutation A (perturb the gated-RMSNorm epsilon inside `GdnBlockPaged`) reds `test_qwen3_5_gdn_spec_routing`'s MIXED case and 5 cases of `test_qwen27_paged_forward`, which is what proves the qwen3.5/3.6 forward still runs this block; mutation B (the wrapper stops delegating) reds only the new case, which is what proves the seam is a wrapper and not a copy. **UNREACHED at this merge commit and deliberately so**, exactly as `RunMoeBlock` landed: no production entry point calls it until the `qwen4_exp` forward lands under [#2031](https://github.com/mudler/vllm.cpp/issues/2031), owned by row `MODEL-MM-QWEN4-EXP` under [#1978](https://github.com/mudler/vllm.cpp/issues/1978). The qwen3.5/3.6 path is byte-identical: the four qwen3_5 suites report the same case and assertion counts before and after (`test_qwen3_5_gdn_spec_routing` 6 cases / 52 assertions on its pre-existing cases, `test_qwen27_paged_forward` 31 / 770, `test_qwen35_moe_gdn_ba_owner` 1 / 23, `test_qwen3_5_decode_graph_seam` 10 / 156) | feature | -| [#2108](https://github.com/mudler/vllm.cpp/issues/2108) | `SPEC-DFLASH2` | **No CI runner has a GPU, so every test that appears to gate a device path is either skipped or silently running on the CPU backend — and both shapes report green.** `.github/workflows/` has no GPU job; `cuda-fat-build` compiles ten architectures and runs nothing. Found while gating SPEC-DFLASH2 W12 D1 ([#2087](https://github.com/mudler/vllm.cpp/issues/2087)), where it bit twice. `tests/vllm/v1/spec_decode/test_dflash_propose.cpp:50` builds its queue with `vt::Queue Cpu()`, so its 10/10 on a GB10 box was a CPU result taken beside an idle GPU and was briefly read as device coverage. And `test_dflash2_runner_reach` is 7-RED under CUDA: measured on `dgx:gpu0` in `vllmcpp-build:gb10` with `--gpus all`, pre-D1 `build18` gives 8 cases / 1 passed / 7 failed / 75 assertions / 18 failed and D1 `build19` gives 9 / 2 / 7 / 83 / 18 — identical failure counts, so PRE-EXISTING and not D1's, and D1's own added case passes. Representative: `:398` `REQUIRE(st_off.block_kernel_calls > 0)` fails as `0 > 0`, a CPU-path route counter that never increments when CUDA is selected; also `:278` and `:345` `REQUIRE_FALSE(blocks.empty())` and `:344`/`:390`/`:391` `CHECK(threw.empty())`. Whether the repair is backend-awareness or a real CUDA-path defect is NOT settled: `:390`/`:391` are `threw` assertions, so something is throwing on the device lane, and a case made to pass by widening its scope would be the failure AGENTS.md names under "Changing the rules or a checker". The consequence for SPEC-DFLASH2 is that `ForwardWithCtxKVDev` at `P > 1` with real device tensors — the path D1 changed — is gated by nothing but an end-to-end throughput run, which an acceptance-only defect is invisible to. Listed under `## Owed` O6 in [`.agents/specs/dflash2-batch-propose.md`](../specs/dflash2-batch-propose.md). | verification | +| [#2108](https://github.com/mudler/vllm.cpp/issues/2108) | `SPEC-DFLASH2` | **No CI runner has a GPU, so every test that appears to gate a device path is either skipped or silently running on the CPU backend — and both shapes report green.** `.github/workflows/` has no GPU job; `cuda-fat-build` compiles ten architectures and runs nothing. Found while gating SPEC-DFLASH2 W12 D1 ([#2087](https://github.com/mudler/vllm.cpp/issues/2087)), where it bit twice. `tests/vllm/v1/spec_decode/test_dflash_propose.cpp:50` builds its queue with `vt::Queue Cpu()`, so its 10/10 on a GB10 box was a CPU result taken beside an idle GPU and was briefly read as device coverage. And `test_dflash2_runner_reach` is 7-RED under CUDA: measured on `dgx:gpu0` in `vllmcpp-build:gb10` with `--gpus all`, pre-D1 `build18` gives 8 cases / 1 passed / 7 failed / 75 assertions / 18 failed and D1 `build19` gives 9 / 2 / 7 / 83 / 18 — identical failure counts, so PRE-EXISTING and not D1's, and D1's own added case passes. Representative: `:398` `REQUIRE(st_off.block_kernel_calls > 0)` fails as `0 > 0`, a CPU-path route counter that never increments when CUDA is selected; also `:278` and `:345` `REQUIRE_FALSE(blocks.empty())` and `:344`/`:390`/`:391` `CHECK(threw.empty())`. Whether the repair is backend-awareness or a real CUDA-path defect is NOT settled: `:390`/`:391` are `threw` assertions, so something is throwing on the device lane, and a case made to pass by widening its scope would be the failure AGENTS.md names under "Changing the rules or a checker". The consequence for SPEC-DFLASH2 is that `ForwardWithCtxKVDev` at `P > 1` with real device tensors — the path D1 changed — is gated by nothing but an end-to-end throughput run, which an acceptance-only defect is invisible to. Listed under `## Owed` O6 in [`.agents/specs/dflash2-batch-propose.md`](specs/dflash2-batch-propose.md). | verification | | [#2064](https://github.com/mudler/vllm.cpp/issues/2064) | `MODEL-MM-QWEN4-EXP` | **`Qwen4ExpHfConfigFromGguf` and `ParseQwen4ExpParams` have never been COMPOSED**, and three keys go missing between them, each silent in a different way. `indexer_kv_heads` is absent and the QSA group is all-or-nothing, so NO real `qwen4exp` file parses at all ("QSA config is missing required fields: indexer_kv_heads"). `ple_layer_ids` is absent and its absence is LEGAL, so the model resolves an EMPTY PLE set — no n-gram table, no PLE layer, `number_of_conv_states()` reporting 1 where the architecture needs 3 — with nothing refusing; W6a left `qwen4exp.ple.layers` unmapped on the ground that neither file says which end the offset is on, and the converter says so in one line (`ple_layers = [i - 1 for i in hp["ple_layer_ids"]]`, llama.cpp #27742 head `035e2273`), so the GGUF key is ZERO-based. `ple_embed_dim` is absent and defaults to `hidden_size`, which is right on the released checkpoint by COINCIDENCE. Invisible to both existing gates because W6a's builds a config and never parses it while W1's parses a config.json and never builds one. Found and FIXED IN FLOW while writing W5a ([#2031](https://github.com/mudler/vllm.cpp/issues/2031)); mutations M11/M12/M13 in `test_qwen4_exp_gguf_weights.cpp` red on each half. | bug | | [#2083](https://github.com/mudler/vllm.cpp/issues/2083) | `MODEL-MM-QWEN4-EXP` | **`qwen4_exp` on `--device cuda` expands the n-gram gather table to 95.4 GiB of anonymous host memory, and the #1123 on-disk device-fit guard structurally cannot see it.** Found by the fresh review of W5a ([#2031](https://github.com/mudler/vllm.cpp/issues/2031)) and fixed in that same flow. `DeviceQuantGatherSupported` (`src/vllm/model_executor/model_loader/gguf_keep_quant.cpp:167`) returns true for `vt::DeviceType::kCPU` and nothing else, because only the CPU `Embedding` kernel decodes a block row. CUDA DOES register `kMatmulBTQuant`, so `GgufLoadPolicy::FromEnv()` sets `keep_quant = true` there and every GEMM weight keeps its blocks — the n-gram GATHER table alone is refused by the device gate in `RouteGgufTensor` and routes to `kExpandBf16`. Derived from the committed 1224-tensor manifest (`tests/vllm/models/qwen4_exp_gguf_manifest.inc`, `unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S` @ `8bdc666649440e9bdc97e16f3f75782c98478ff5`): `per_layer_token_embd.weight` is [320001536, 160] IQ4_NL, **28,800,138,240 B = 26.822 GiB on disk** and **320001536 x 160 x 2 = 102,400,491,520 B = 95.368 GiB expanded**, on a box with ~119.6 GiB for everything. The #1123 guard (`src/vllm/entrypoints/model_loader.cpp:2389`) sums the file's ON-DISK tensor bytes, which are **72,535,436,800 B = 67.554 GiB** and comfortably inside the budget, so it admits the load and the expansion happens after it — `model_loader.cpp`'s own stated worst case, "Loading for 26 minutes and dying mid-stream is the worst of the available behaviours". **The second half of the issue is a behaviour regression W5a introduced and its own record did not name:** before W5a, pointing the loader at a `qwen4exp` GGUF refused immediately because `load_weights` was a refuse-by-name stub; after W5a it performs the full load and then dies in `MakeQwen4ExpKVCache`, which still refuses because W5c owes the KV-cache spec. **Fixed in flow:** `LoadQwen4ExpFromGguf` now takes the device the forward will run on — NO default, so a caller cannot disable the guard by saying nothing — and refuses BY NAME ahead of any tensor I/O when the config names a PLE layer and `DeviceQuantGatherSupported(device)` is false, naming the tensor, the device, the missing gather arm and `--device cpu`. Gated red-first by `tests/vllm/models/test_qwen4_exp_gguf_weights.cpp` "a device with no block gather refuses BEFORE the load": 8 assertions red before the guard, green after, and mutation MUT-G1 (`if (false && ...)`) reddens the same 8. **NOT closed by that fix and recorded under the row's `## Owed`:** the CUDA block-decoding gather kernel itself, so this stays a CPU-only arm; and the load-then-die ordering is only NARROWED, since on `--device cpu` the load still completes before `MakeQwen4ExpKVCache` throws — W5c closes that. **One honest gap in the gate, stated rather than implied:** mutation MUT-G2, pinning the production device ARGUMENT in `qwen4_exp_registry.cpp` to a literal `vt::DeviceType::kCPU`, SURVIVES and cannot do otherwise on a CPU-only build, where no other platform is registered and `CurrentPlatform().device_type()` and the literal are the same value. The guard is gated; the argument is not; closing that needs a CUDA host | bug | -| [#2081](https://github.com/mudler/vllm.cpp/issues/2081) | `MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation` | **Nothing in the tree can tell the shipped `qwen3_5` GGUF V-head reorder from its own inverse.** `ReorderVRows`/`ReorderVCols` (`src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp`) recover HF grouped order from GGUF tiled order by reading grouped head `g = k*R + r` out of tiled head `t = r*K + k`. Mutation MUT-M6 swaps the two definitions, which applies the map in the wrong direction. Re-measured on `row/MODEL-MM-QWEN4-EXP-W5` rather than relayed: it **SURVIVES all four suites**, each byte-for-byte the un-mutated count -- `test_gguf_qwen36_loader` 7/7 cases, 555 assertions; `test_model_loader_gguf` 7/7, 23; `test_gguf_nvfp4` 14/14, 2352; `test_gguf_keep_quant` 42/42, 6340. The cause is the fixtures, not the loader: every synthetic `qwen35`/`qwen35moe` GGUF in the tree states `ssm.group_count = 2` with `ssm.time_step_rank` 2 or 4, so `K = 2` and `R` is 1 or 2. The ONE case that exercises the reorder at all, `test_gguf_qwen36_loader.cpp`'s "V-head reorder when num_v != num_k", is `K == R == 2`, where the permutation is its own INVERSE and the mutated loader emits byte-identical weights; every other fixture is `R = 1`, where it is the identity. The 27B manifest is the only `K != R` shape in the tree (`group_count 16`, `time_step_rank 48`) and it carries metadata only, so no buffer passes through the reorder there. This is the loader that SHIPS for `qwen35`, `qwen35moe` and `qwen3next`, so the untested direction is a live silent-wrong-weights risk on real files. Found by the fresh review of W5a of [#2031](https://github.com/mudler/vllm.cpp/issues/2031), which deliberately duplicates the function for `qwen4_exp`; OUR copy is gated -- its fixture is `K = 2, R = 3` and the same swap REDs 2 of 11 cases and 41 assertions (mutation M5). **Not fixed in that flow**, because closing it re-shapes a shipped model's fixtures and moves `qwen35`, `qwen35moe` and `qwen3next` coverage, which is outside `MODEL-MM-QWEN4-EXP`'s scope. Listed under `## Owed` in [`specs/qwen4-exp-flash-next.md`](../specs/qwen4-exp-flash-next.md) until the owning row picks it up | gap | -| [#2111](https://github.com/mudler/vllm.cpp/issues/2111) | `SPEC-DFLASH2` | **W12 D2's two terms are bounded at ~2% of the c=8 step — under the wave's own ~6% resolution floor — and the 1.78x it is aimed at is a scaling no engine on this box reaches with speculation ON.** D2 (`## Design` of [`specs/dflash2-batch-propose.md`](../specs/dflash2-batch-propose.md)) is one shared paged pool plus a batched block table, so the batched propose becomes the `P == 1` path at `num_reqs > 1`; it is feasible and its change surface is fully traced in the spec, and it is NOT being implemented. Post-D1 the `P > 1` lane's whole `O(C)` cost is two `IndexSelect`/`IndexCopy` stages — the gather in `ForwardBlockLogitsWithDeviceKV` (16 B per context row per `kdim` element per layer) and the combined scatter in `ForwardWithCtxKVDev` (8 B on the same basis) — so `B_saved = 24 * L * C * kdim`. With `L = 5` (every run reports `dflash=5`), `C ~ 8 x 1300` at c=8, and `kdim` BOUNDED rather than read because O3 is still owed (`kdim <= hidden_size = 5120`, realistically 512-1024 under GQA), that is 639 MB / **2.3 ms** at GB10's ~273 GB/s, and 23.4 ms even at the impossible `Hkv == Hq` ceiling. The graph lane D2 also restores removes host dispatch, not device work: ~255 launches at `P=8, L=5` puts its ENTIRE ceiling at ~1.5 ms. Against a 203 ms c=8 step (`76.23 / (8 * 1.9342)`), reaching 112 tok/s needs 64.8 ms removed and reaching SGLang's scaling needs 29.9 ms; D2's two terms are ~3.8 ms, or 5.9% and 12.7% of those. **The premise is also wrong, and that correction is the more useful half:** vLLM scales 1.245x from c=4 to c=8 with speculation on, SGLang 1.418x, ours 1.208x, against speculation-OFF's 1.777x — no engine reaches 1.78 with speculation on, because at `k = 8` and c=8 the verify batch is 72 target rows against speculation-off's 8, so a speculative step does ~9x the target work by construction while the ladder counts accepted tokens. Reading 1.208 against 1.777 charges the draft for the verify batch's growth; the reachable denominator is SGLang's 1.418. **Where the residual is remains unattributed:** the non-speculative step grows +12.9 ms across that rung and the speculative step +80.4 ms, split between the verify batch and the draft phase, and nothing has measured the split — E1 (`VT_SPEC_TRACE=2`, `[spec-phase-dev] pre= fwd= select= walk=`) prints exactly it, needs no code change, and O5 records it was traced only at c=1; E6 is likewise unread and free. D2's own precondition is unmet as well: `## Design` conditions the pool on agreeing the allocation with [#2007](https://github.com/mudler/vllm.cpp/issues/2007), which is open and unowned, and the arena is a residency change from today's lazy `first_sight` per-request allocation to `max_num_reqs` up front — up to the whole 8 GiB `kDflashCtxTotalBudgetBytes` aggregate — on the box #1647 OOM-rebooted. NO new device measurement: every figure is arithmetic over code that was read and over the ladders the spec already records, acceptance is measured only at c=8 (1.9342, n=532) and assumed equal at c=4, and every rung is n=1 | perf | -| [#2112](https://github.com/mudler/vllm.cpp/issues/2112) | `SPEC-DFLASH2` | **E6 cannot be run and Gate 3 cannot be read on a server: `GraphDispatchStats` and `DflashBlockRouteStats` have no production reader, so both counter families are visible only to tests.** Found while bounding W12 D2 ([#2111](https://github.com/mudler/vllm.cpp/issues/2111)). Every caller of `vllm::v1::GetGraphDispatchStats()` and `vllm::detail::GetDflashBlockRouteStats()` is under `tests/`; `grep -rn "graph-dispatch\|GRAPH_STATS" src/vllm/v1/worker/gpu/runner.cpp` is empty, and the counters are process-local diagnostics the server never surfaces. **E6 of [`specs/dflash2-batch-propose.md`](../specs/dflash2-batch-propose.md) is therefore not runnable as written** — it says to read `GraphDispatchStats::uniform_spec_steps / total` and `spec_as_decode_steps` at c=4 and c=8, ranks itself in the first group to run ("Run E1, E2 and E6 first"), calls the whole-step raggedness cliff it measures a mechanism INDEPENDENT of that wave, and says it "costs nothing to read alongside them"; it costs a code change. **And Gate 3 cannot be read on the ladder either:** [#2089](https://github.com/mudler/vllm.cpp/issues/2089) landed `DflashBlockRouteStats` with W12 D1 so a gate could assert which attention lane a draft block took, and D1 added `NoteDflashCombinedAttn` for the `P > 1` lane the counters had been blind to — correct in the forward, complete in the tests, and unreadable on the c=1/2/4/8 ladder, which is the only place the `P > 1` lane runs in anger and the reason the counter exists. **This is the #2089 shape one level out.** That wave's lesson was "an instrument that only counts the fast path cannot report that a slow path exists"; the repair widened the counter to both lanes but gave neither a READOUT outside a unit test, so the instrument is complete and still reports nothing on the workload it was built for — a counter whose only reader is a test measures a class, not a capability, which is [`reachability.md`](../reachability.md)'s distinction applied to a diagnostic. Scope: an env-gated line where `[spec-phase]` / `[spec-phase-dev]` already print from (`src/vllm/v1/worker/gpu/runner.cpp`), reporting both families, host-only and CPU-gateable. NOT fixed in the flow that found it, which is a record-only refutation of D2 carrying no product code; listed under `## Owed` O7 in that spec | verification | -| [#2123](https://github.com/mudler/vllm.cpp/issues/2123) | `MODEL-MM-QWEN4-EXP` | **The `qwen4_exp` gated-residual hyper-connection stream had no device arm, and the shared `vt::` surface cannot express one.** `Qwen4ExpTextDecoderLayer` reads it twice per layer over 48 layers and `Qwen4ExpTextModel` once more for the terminal `use_combine=False` mixer — 97 sites — so it is the structure the whole forward is threaded through, and W3 ([#1988](https://github.com/mudler/vllm.cpp/issues/1988)) had landed the arithmetic only as a `std::vector` HOST reference, which the row's own `## Owed` records as an oracle for a device arm and not a substitute for one. The gap is measured rather than asserted, surveyed at `331eda888`: no ungated per-group RMS norm (`vt::RmsNormGated` has no `group_size`; `vt::RmsNormGatedGroup` takes a non-nullable SILU gate), `vt::RmsNorm` cannot carry a per-group weight (one `[H]` gamma per row against `hc_norm`'s `[hc*H]`), and no standalone `silu`, `sigmoid`, elementwise binary multiply or axis reduction exists anywhere in the op set — `kSilu`/`kSigmoid`/`kMul` are `FOp` opcodes inside a `constexpr FusedRecipe` and are unreachable as free functions. A composition would have needed five new general ops and would still have materialised the `[T, hc, H]` broadcast the rank-1 write-back exists to avoid, which is what both llama.cpp implementations of this architecture do. Closed by two new `vt::` ops on the `kDeepseekV4Mhc` precedent — `Qwen4ExpGatedResidual` (grouped norm, low-rank SiLU gate, sigmoid, mean collapse, injection weights) and `Qwen4ExpGatedResidualWriteBack` (the rank-1 update, in place) — batched over T tokens, CPU kernel in `src/vt/cpu/cpu_qwen4_exp.cpp`, gated against the SAME lane-pinned transformers v5.16.0 goldens the host reference answers to plus a model-width (2560/4/320) agreement check the golden shapes cannot make. 18-mutation battery, 16 RED: the two survivors of the first battery were REPAIRED rather than recorded, and both repairs closed a hole that pre-dated this wave — the epsilon placement (`+ eps` inside the rsqrt) was invisible at every existing golden's magnitude and survived in the W3 HOST suite too, closed by a fourth golden case at `hyper_scale = 0.01` generated from the same pinned oracle and driven by both suites (M7 device / M16 host now red, A/B/C byte-identical); and the `double` per-group accumulator, closed by a zeroed-projection case at group size 2560 where a `float` accumulator is 571x over the bound. One survivor stands and is named: M14, `x / hc` versus `x * (1.0f / hc)`, a sub-ulp difference no f32 golden comparison can separate. **UNREACHED at this merge commit and deliberately so**, exactly as W2/W3/W4 landed: the architecture's only production entry point is `ModelRegistry::Forward`, which is all-or-nothing, so every slice below the whole forward is unreached by construction; the wiring is owed by W5b under [#2031](https://github.com/mudler/vllm.cpp/issues/2031) and tracked by [#1978](https://github.com/mudler/vllm.cpp/issues/1978). Also owed and named in the spec's `## Owed`: the CUDA arm (unwritten because it could not be gated on this CPU host) and the reduction width it has to choose, and one deliberate divergence from upstream's `.type_as(x)` intermediate rounding in the bf16 arm | feature | -| [#2131](https://github.com/mudler/vllm.cpp/issues/2131) | `MODEL-MM-QWEN4-EXP` | **The GPU runner accepts exactly ONE `MambaSpec` group of exactly TWO shapes, so a PLE layer's third conv state is unrepresentable and W5c's KV-cache spec cannot be written against it.** Verified in the tree rather than relayed: `src/vllm/v1/worker/gpu/runner.cpp` asserts `mamba_spec->shapes.size() == 2 && mamba_spec->dtypes.size() == 2` ("runner: recurrent MambaSpec must contain conv then temporal state") and then reads `shapes[0]` as conv and `shapes[1]` as temporal, while the topology refusal a few lines above states the one-group rule in prose ("any number of non-eagle AttentionSpec groups and ONE MambaSpec group") and `gdn_group_id_` is a single scalar index, not a list. A `qwen4_exp` PLE layer carries THREE persistent recurrent streams — the GDN conv, the PLE dilated depthwise conv (`kernel_size = 4, dilation = 3`, a 9-deep ring buffer read at stride 3) and the int64 n-gram token history — so the two-shape assumption cannot express the states and the one-group assumption cannot address them separately. The blocker is in the ENGINE, not in the model registry, which is why it is its own issue rather than part of a model wave. Scope: generalise the recurrent-cache topology so a `MambaSpec` group can carry more than two shapes, or so more than one such group can exist, with per-state addressing, mirroring vLLM where it defines the behaviour; the existing Mamba/GDN arms must stay byte-identical, with `test_qwen27_paged_forward` and the qwen3.5 recurrent suites as the regression gate. Recorded under `## Owed` and `## Now` in [`specs/qwen4-exp-flash-next.md`](../specs/qwen4-exp-flash-next.md), which until now described the constraint in prose with no issue behind it | gap | +| [#2081](https://github.com/mudler/vllm.cpp/issues/2081) | `MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation` | **Nothing in the tree can tell the shipped `qwen3_5` GGUF V-head reorder from its own inverse.** `ReorderVRows`/`ReorderVCols` (`src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp`) recover HF grouped order from GGUF tiled order by reading grouped head `g = k*R + r` out of tiled head `t = r*K + k`. Mutation MUT-M6 swaps the two definitions, which applies the map in the wrong direction. Re-measured on `row/MODEL-MM-QWEN4-EXP-W5` rather than relayed: it **SURVIVES all four suites**, each byte-for-byte the un-mutated count -- `test_gguf_qwen36_loader` 7/7 cases, 555 assertions; `test_model_loader_gguf` 7/7, 23; `test_gguf_nvfp4` 14/14, 2352; `test_gguf_keep_quant` 42/42, 6340. The cause is the fixtures, not the loader: every synthetic `qwen35`/`qwen35moe` GGUF in the tree states `ssm.group_count = 2` with `ssm.time_step_rank` 2 or 4, so `K = 2` and `R` is 1 or 2. The ONE case that exercises the reorder at all, `test_gguf_qwen36_loader.cpp`'s "V-head reorder when num_v != num_k", is `K == R == 2`, where the permutation is its own INVERSE and the mutated loader emits byte-identical weights; every other fixture is `R = 1`, where it is the identity. The 27B manifest is the only `K != R` shape in the tree (`group_count 16`, `time_step_rank 48`) and it carries metadata only, so no buffer passes through the reorder there. This is the loader that SHIPS for `qwen35`, `qwen35moe` and `qwen3next`, so the untested direction is a live silent-wrong-weights risk on real files. Found by the fresh review of W5a of [#2031](https://github.com/mudler/vllm.cpp/issues/2031), which deliberately duplicates the function for `qwen4_exp`; OUR copy is gated -- its fixture is `K = 2, R = 3` and the same swap REDs 2 of 11 cases and 41 assertions (mutation M5). **Not fixed in that flow**, because closing it re-shapes a shipped model's fixtures and moves `qwen35`, `qwen35moe` and `qwen3next` coverage, which is outside `MODEL-MM-QWEN4-EXP`'s scope. Listed under `## Owed` in [`specs/qwen4-exp-flash-next.md`](specs/qwen4-exp-flash-next.md) until the owning row picks it up | gap | +| [#2111](https://github.com/mudler/vllm.cpp/issues/2111) | `SPEC-DFLASH2` | **W12 D2's two terms are bounded at ~2% of the c=8 step — under the wave's own ~6% resolution floor — and the 1.78x it is aimed at is a scaling no engine on this box reaches with speculation ON.** D2 (`## Design` of [`specs/dflash2-batch-propose.md`](specs/dflash2-batch-propose.md)) is one shared paged pool plus a batched block table, so the batched propose becomes the `P == 1` path at `num_reqs > 1`; it is feasible and its change surface is fully traced in the spec, and it is NOT being implemented. Post-D1 the `P > 1` lane's whole `O(C)` cost is two `IndexSelect`/`IndexCopy` stages — the gather in `ForwardBlockLogitsWithDeviceKV` (16 B per context row per `kdim` element per layer) and the combined scatter in `ForwardWithCtxKVDev` (8 B on the same basis) — so `B_saved = 24 * L * C * kdim`. With `L = 5` (every run reports `dflash=5`), `C ~ 8 x 1300` at c=8, and `kdim` BOUNDED rather than read because O3 is still owed (`kdim <= hidden_size = 5120`, realistically 512-1024 under GQA), that is 639 MB / **2.3 ms** at GB10's ~273 GB/s, and 23.4 ms even at the impossible `Hkv == Hq` ceiling. The graph lane D2 also restores removes host dispatch, not device work: ~255 launches at `P=8, L=5` puts its ENTIRE ceiling at ~1.5 ms. Against a 203 ms c=8 step (`76.23 / (8 * 1.9342)`), reaching 112 tok/s needs 64.8 ms removed and reaching SGLang's scaling needs 29.9 ms; D2's two terms are ~3.8 ms, or 5.9% and 12.7% of those. **The premise is also wrong, and that correction is the more useful half:** vLLM scales 1.245x from c=4 to c=8 with speculation on, SGLang 1.418x, ours 1.208x, against speculation-OFF's 1.777x — no engine reaches 1.78 with speculation on, because at `k = 8` and c=8 the verify batch is 72 target rows against speculation-off's 8, so a speculative step does ~9x the target work by construction while the ladder counts accepted tokens. Reading 1.208 against 1.777 charges the draft for the verify batch's growth; the reachable denominator is SGLang's 1.418. **Where the residual is remains unattributed:** the non-speculative step grows +12.9 ms across that rung and the speculative step +80.4 ms, split between the verify batch and the draft phase, and nothing has measured the split — E1 (`VT_SPEC_TRACE=2`, `[spec-phase-dev] pre= fwd= select= walk=`) prints exactly it, needs no code change, and O5 records it was traced only at c=1; E6 is likewise unread and free. D2's own precondition is unmet as well: `## Design` conditions the pool on agreeing the allocation with [#2007](https://github.com/mudler/vllm.cpp/issues/2007), which is open and unowned, and the arena is a residency change from today's lazy `first_sight` per-request allocation to `max_num_reqs` up front — up to the whole 8 GiB `kDflashCtxTotalBudgetBytes` aggregate — on the box #1647 OOM-rebooted. NO new device measurement: every figure is arithmetic over code that was read and over the ladders the spec already records, acceptance is measured only at c=8 (1.9342, n=532) and assumed equal at c=4, and every rung is n=1 | perf | +| [#2112](https://github.com/mudler/vllm.cpp/issues/2112) | `SPEC-DFLASH2` | **E6 cannot be run and Gate 3 cannot be read on a server: `GraphDispatchStats` and `DflashBlockRouteStats` have no production reader, so both counter families are visible only to tests.** Found while bounding W12 D2 ([#2111](https://github.com/mudler/vllm.cpp/issues/2111)). Every caller of `vllm::v1::GetGraphDispatchStats()` and `vllm::detail::GetDflashBlockRouteStats()` is under `tests/`; `grep -rn "graph-dispatch\|GRAPH_STATS" src/vllm/v1/worker/gpu/runner.cpp` is empty, and the counters are process-local diagnostics the server never surfaces. **E6 of [`specs/dflash2-batch-propose.md`](specs/dflash2-batch-propose.md) is therefore not runnable as written** — it says to read `GraphDispatchStats::uniform_spec_steps / total` and `spec_as_decode_steps` at c=4 and c=8, ranks itself in the first group to run ("Run E1, E2 and E6 first"), calls the whole-step raggedness cliff it measures a mechanism INDEPENDENT of that wave, and says it "costs nothing to read alongside them"; it costs a code change. **And Gate 3 cannot be read on the ladder either:** [#2089](https://github.com/mudler/vllm.cpp/issues/2089) landed `DflashBlockRouteStats` with W12 D1 so a gate could assert which attention lane a draft block took, and D1 added `NoteDflashCombinedAttn` for the `P > 1` lane the counters had been blind to — correct in the forward, complete in the tests, and unreadable on the c=1/2/4/8 ladder, which is the only place the `P > 1` lane runs in anger and the reason the counter exists. **This is the #2089 shape one level out.** That wave's lesson was "an instrument that only counts the fast path cannot report that a slow path exists"; the repair widened the counter to both lanes but gave neither a READOUT outside a unit test, so the instrument is complete and still reports nothing on the workload it was built for — a counter whose only reader is a test measures a class, not a capability, which is [`reachability.md`](reachability.md)'s distinction applied to a diagnostic. Scope: an env-gated line where `[spec-phase]` / `[spec-phase-dev]` already print from (`src/vllm/v1/worker/gpu/runner.cpp`), reporting both families, host-only and CPU-gateable. NOT fixed in the flow that found it, which is a record-only refutation of D2 carrying no product code; listed under `## Owed` O7 in that spec | verification | | [#2133](https://github.com/mudler/vllm.cpp/issues/2133) | `ENG-HYBRID-PLACEMENT` | Kimi-Linear WAS wireable and the seam's record said otherwise: its `MoeBlockDevice`/`MoeBlockDeviceBf16` are seam-shaped, but a sweep read `kimi_linear_forward.cpp`, saw a host `std::vector` path, and generalised one file to the whole architecture — the same error already made about Laguna. Five architectures now route through the seam, and the remaining four are recorded by REASON rather than as one undifferentiated list | bug | -| [#2117](https://github.com/mudler/vllm.cpp/issues/2117) | `SPEC-DFLASH2` | **Two silent graph-loss mechanisms on the c=8 spec step, and the cost model that made the first look measurable.** A step that admits a prefill loses the CUDA graph for the WHOLE step (`GraphEligibleQueryLen` returns `nullopt`, `cudagraph_dispatch.h:161-176`) AND drops its decode rows onto the `num_splits=1` prefill ladder (`uniform_spec_query_len` becomes 0, `runner.cpp:2065-2069`, and `PagedAttnIsPrefill` then reads true, `include/vt/paged_attn_route.h:47-51`). W13 lands the INSTRUMENT and corrects the arithmetic rather than fixing either mechanism: #2117 prices mechanism 1 by multiplying the mixed step's TOTAL cost by its frequency, but ~1976 of a 2048-token mixed step is prefill work the engine owes on any lane, so the defect's MARGINAL cost is only the eager term (~24% of a decode step at c=64, an upper bound at c=8, `runner.cpp:2003-2004`) plus the attention-lane term (+9 ms on a ~113-122 ms step, `specs/dflash2-spec-as-decode.md`), which over 3% to 7% of steps is **1.0% to 2.2%, under the rung's own 5.9% spread** — mechanism 1 can be neither confirmed nor refuted by a c=8 throughput delta, so a fix landed against it now would be unfalsifiable in both directions. What ships instead is the readout #2112 owes plus a THREE-WAY split of `ragged_steps` into `ragged_mixed_steps`, `ragged_prefill_only_steps` and `ragged_spec_only_steps`, because the flat counter is consistent both with #2117's 3-7% admission prediction and with its own "far above 10% means #1943", and therefore discriminates nothing. All three candidate fixes are assessed and NONE lands: **(a) PIECEWISE** is a real primary-oracle gap (`cudagraph_dispatcher.py:307-324`, ours goes fully eager, `cudagraph_dispatch.h:204-207` already names it owed) but `include/vt/breakable_graph.h:21-26` states the seam's own admissibility rule — a speed claim needs a path "BOTH currently eager AND currently host-bound" — and a 2048-token mixed step is the prefill regime, 3.8% host-idle at >96% GPU-busy, so the lever is inert there and the construction is several waves without a compiler; **(b) keep the decode rows on the decode lane** is the half with a measured number, and its unlanded PRECONDITION is itself a mirror gap this issue found — `runner.cpp:1691` calls `reorder_batch_to_split_decodes_and_prefills` with the default `decode_threshold = 1` while vLLM passes `reorder_batch_threshold`, which `_init_reorder_batch_threshold` (`backend.py:657-687` @ pin `5559679229`) raises to `1 + 2k` for a spec-as-decode backend, so at k=8 the threshold should be 17 and is 1 and a 9-token verify row sorts as `long_extend` beside chunked-prefill continuations, leaving no boundary any split could cut at; **(c) SGLang's prefill/decode step separation** is a divergence from the mirror source, not a porting gap, and is REJECTED pending the measurement. NOT re-run, and named so nobody re-runs it: the `mnbt=2048` vs `8192` A/B at `model_loader.cpp:1100-1103`, whose stated reason is mechanism 1 in its own words. No GPU taken, no throughput claim made; O1/O2/O3 in [`specs/dflash2-mixed-step-readout.md`](../specs/dflash2-mixed-step-readout.md) carry the remainder | bug | -| [#2129](https://github.com/mudler/vllm.cpp/issues/2129) | `SPEC-DFLASH2` | **The spec-as-decode reorder threshold is computed and never passed to the reorder, so the runner reorders every batch at 1.** `runner.cpp:1691` calls `reorder_batch_to_split_decodes_and_prefills(input_batch_, scheduler_output)` with no `decode_threshold`, taking the declaration default of 1 (`runner.h:102-104`), while `SpecAsDecodeReorderThreshold` (`include/vllm/v1/attention/backend.h:180-184`) already mirrors upstream's `1 + (parallel_drafting ? 2 : 1) * k` and has exactly ONE caller — `backend.h:201`, inside the same header, serving the spec-as-decode *classification* that `runner.cpp:2065` consumes. Upstream feeds it: `_may_reorder_batch` passes `decode_threshold=self.reorder_batch_threshold` (`gpu_model_runner.py:1126-1130` @ pin `5559679229`), `calculate_reorder_batch_threshold` resolves it as the min over the attention groups after their builders exist (`:7122`, `:7194-7212`), and `_init_reorder_batch_threshold` raises it for a spec-as-decode backend (`backend.py:657-687`), which `gdn_attn.py:112` requests for every speculative configuration via `supports_spec_as_decode=self.use_spec_decode`. At k=8 with `parallel_drafting` — true for `dflash` and `dspark` and only those (`speculative.py:963-964`, mirrored at `include/vllm/config/speculative.h:158`, `:346`) — the threshold should be 17 and is 1, so a `1+k = 9`-token verify row classifies `long_extend` (region 2) instead of `decode` (region 0) and sorts among the chunked-prefill continuations, leaving no decode/prefill boundary for any consumer of the reordered order to cut at. **A mirror obligation, not a speed claim**: the formula was ported by W10 and left unwired, and NOTHING here is measured — the c=8 floor on the #1574 ladder is ~5.9 %. Found by the #2117 implementer while scoping the mixed-step work and named in [#2128](https://github.com/mudler/vllm.cpp/issues/2128) as lever (b)'s unlanded precondition; verified independently before dispatch. Fixed in flow, spec [reorder-threshold-wiring.md](../specs/reorder-threshold-wiring.md), which also records under `## Owed` the divergence this change does NOT close: upstream skips the reorder entirely when every group reports `None` (`flash_attn.py` and `triton_attn.py` set no threshold), and we reorder unconditionally | bug | -| [#2116](https://github.com/mudler/vllm.cpp/issues/2116) | `SPEC-DFLASH2` | **A speculator vetoes the async input and sampler path at `runner.cpp:470` (`:425` at the base tree the measurement below names), so every spec step drains the queue in step, while vLLM keeps async scheduling ON for dflash because `DFlashModelTypes` is inside `EagleModelTypes`.** Scoped by [`specs/dflash2-async-spec-sampler.md`](../specs/dflash2-async-spec-sampler.md), which discharges `## Owed` A2 of [`specs/spec-decode-async-scheduling.md`](../specs/spec-decode-async-scheduling.md). The veto was MEASURED load-bearing on the CPU tier rather than argued: deleting `!spec_config_.has_value()` at both construction sites reds `test_mtp_depth`'s W7 identity case (10 cases / 123 assertions / exit 0 becomes 9 passed / 1 failed / exit 1) through a production refusal at `runner.cpp:1833`, because `sample_tokens_async` carries no verify arm — no rejection sampler and no propose — a reason the veto's own comment did not name. The comment's stated reason holds too, and holds invisibly: under the same mutation the non-draft-aware combine overwrites the LAST DRAFT of every verify block with the previous step's committed token (`draft=[6 18]` becomes `draft=[6 5]` where the previous step emitted `5`, at every position), and the emitted tokens never move, so every identity assertion still passes. That is #1366's acceptance-only shape a second time. The row therefore stays vetoed and the fix is staged A2-1 through A2-5, with the draft-equality gate G2 owed by the first wave | bug | +| [#2117](https://github.com/mudler/vllm.cpp/issues/2117) | `SPEC-DFLASH2` | **Two silent graph-loss mechanisms on the c=8 spec step, and the cost model that made the first look measurable.** A step that admits a prefill loses the CUDA graph for the WHOLE step (`GraphEligibleQueryLen` returns `nullopt`, `cudagraph_dispatch.h:161-176`) AND drops its decode rows onto the `num_splits=1` prefill ladder (`uniform_spec_query_len` becomes 0, `runner.cpp:2065-2069`, and `PagedAttnIsPrefill` then reads true, `include/vt/paged_attn_route.h:47-51`). W13 lands the INSTRUMENT and corrects the arithmetic rather than fixing either mechanism: #2117 prices mechanism 1 by multiplying the mixed step's TOTAL cost by its frequency, but ~1976 of a 2048-token mixed step is prefill work the engine owes on any lane, so the defect's MARGINAL cost is only the eager term (~24% of a decode step at c=64, an upper bound at c=8, `runner.cpp:2003-2004`) plus the attention-lane term (+9 ms on a ~113-122 ms step, `specs/dflash2-spec-as-decode.md`), which over 3% to 7% of steps is **1.0% to 2.2%, under the rung's own 5.9% spread** — mechanism 1 can be neither confirmed nor refuted by a c=8 throughput delta, so a fix landed against it now would be unfalsifiable in both directions. What ships instead is the readout #2112 owes plus a THREE-WAY split of `ragged_steps` into `ragged_mixed_steps`, `ragged_prefill_only_steps` and `ragged_spec_only_steps`, because the flat counter is consistent both with #2117's 3-7% admission prediction and with its own "far above 10% means #1943", and therefore discriminates nothing. All three candidate fixes are assessed and NONE lands: **(a) PIECEWISE** is a real primary-oracle gap (`cudagraph_dispatcher.py:307-324`, ours goes fully eager, `cudagraph_dispatch.h:204-207` already names it owed) but `include/vt/breakable_graph.h:21-26` states the seam's own admissibility rule — a speed claim needs a path "BOTH currently eager AND currently host-bound" — and a 2048-token mixed step is the prefill regime, 3.8% host-idle at >96% GPU-busy, so the lever is inert there and the construction is several waves without a compiler; **(b) keep the decode rows on the decode lane** is the half with a measured number, and its unlanded PRECONDITION is itself a mirror gap this issue found — `runner.cpp:1691` calls `reorder_batch_to_split_decodes_and_prefills` with the default `decode_threshold = 1` while vLLM passes `reorder_batch_threshold`, which `_init_reorder_batch_threshold` (`backend.py:657-687` @ pin `5559679229`) raises to `1 + 2k` for a spec-as-decode backend, so at k=8 the threshold should be 17 and is 1 and a 9-token verify row sorts as `long_extend` beside chunked-prefill continuations, leaving no boundary any split could cut at; **(c) SGLang's prefill/decode step separation** is a divergence from the mirror source, not a porting gap, and is REJECTED pending the measurement. NOT re-run, and named so nobody re-runs it: the `mnbt=2048` vs `8192` A/B at `model_loader.cpp:1100-1103`, whose stated reason is mechanism 1 in its own words. No GPU taken, no throughput claim made; O1/O2/O3 in [`specs/dflash2-mixed-step-readout.md`](specs/dflash2-mixed-step-readout.md) carry the remainder | bug | +| [#2116](https://github.com/mudler/vllm.cpp/issues/2116) | `SPEC-DFLASH2` | **A speculator vetoes the async input and sampler path at `runner.cpp:470` (`:425` at the base tree the measurement below names), so every spec step drains the queue in step, while vLLM keeps async scheduling ON for dflash because `DFlashModelTypes` is inside `EagleModelTypes`.** Scoped by [`specs/dflash2-async-spec-sampler.md`](specs/dflash2-async-spec-sampler.md), which discharges `## Owed` A2 of [`specs/spec-decode-async-scheduling.md`](specs/spec-decode-async-scheduling.md). The veto was MEASURED load-bearing on the CPU tier rather than argued: deleting `!spec_config_.has_value()` at both construction sites reds `test_mtp_depth`'s W7 identity case (10 cases / 123 assertions / exit 0 becomes 9 passed / 1 failed / exit 1) through a production refusal at `runner.cpp:1833`, because `sample_tokens_async` carries no verify arm — no rejection sampler and no propose — a reason the veto's own comment did not name. The comment's stated reason holds too, and holds invisibly: under the same mutation the non-draft-aware combine overwrites the LAST DRAFT of every verify block with the previous step's committed token (`draft=[6 18]` becomes `draft=[6 5]` where the previous step emitted `5`, at every position), and the emitted tokens never move, so every identity assertion still passes. That is #1366's acceptance-only shape a second time. The row therefore stays vetoed and the fix is staged A2-1 through A2-5, with the draft-equality gate G2 owed by the first wave | bug | | [#2141](https://github.com/mudler/vllm.cpp/issues/2141) | `MODEL-MM-QWEN4-EXP` | **Supersedes one claim in the [#2123](https://github.com/mudler/vllm.cpp/issues/2123) row above: the epsilon-placement defect did NOT survive in the W3 HOST suite.** That row says it did. W3 gated the placement deliberately, at `tests/vllm/models/test_qwen4_exp_hc.cpp:268-276`, inside "qwen4_exp grouped RMSNorm mirrors RMSNormGated(group_size) at the lane pin" — a `big_eps = 4.0f` probe against the double reference `NormRefD`, whose own comment says an eps-placement defect is invisible to every golden at the model's real `1e-6` and that "a case at an eps large enough to separate them is the only thing that gates it". Pre-repair, with `+ eps` moved outside the rsqrt, that probe is RED at `CHECK( 0.802185 < 1e-05 )` — 2 of 14 cases at `origin/main` — while the unmutated kernel clears the same `kTol = 1e-5` by `7.77e-08`, so it discriminates by seven orders of magnitude. The ungated arm was the DEVICE suite ALONE, and W5b-2's golden case D at `hyper_scale = 0.01` is what shuts it. Case D additionally sharpens the HOST arm from those 2 red assertions to 10 (M16: RED, 3 of 15 cases), which is an ENHANCEMENT of a gate that already fired and not a hole closed. The [#2123](https://github.com/mudler/vllm.cpp/issues/2123) row is not edited and never will be: `merge=union` DUPLICATES an edited row instead of merging it, so the immutability protects the merge driver rather than the error, and an appended superseding row is what an append-only log is for. A second row keyed on #2123 is equally impossible — `check-agent-record.py` reports a repeated issue number as the duplicate two branches appending the same issue would produce — hence this row's own issue. The spec, its `## Mutation record — W5b-2`, the M16 table row and the case-D comment in the test file were already corrected in `11a61e3bd` | record | | [#2077](https://github.com/mudler/vllm.cpp/issues/2077) | `SERVE-CHAT-TEMPLATE` | server: GGUF models fall back to the naive role-join prompt because the chat template is never loaded from GGUF metadata — `server_main.cpp` calls `LoadChatTemplateFromConfig` only, and a `.gguf` file has no `tokenizer_config.json`, so the model loops on `\nassistant:\n` with zero useful output | bug | -| [#2119](https://github.com/mudler/vllm.cpp/issues/2119) | `SAMPLE-N-ASYNC` | **`/v1/completions` streaming drops every choice past the first**: `CompletionSseStream::next` (`src/vllm/entrypoints/openai/serving_completion.cpp:92`) reads `response.outputs.front()` and formats one SSE choice from it, while `RequestOutputCollector::Merge` (`src/vllm/v1/engine/output_processor.cpp:96-118`) keeps distinct `index` completions side by side in ONE frame whenever the producer outruns the consumer. Upstream flattens the samples one choice per chunk and asserts it (`tests/entrypoints/openai/completion/test_completion.py:419-424`, `test_parallel_streaming` at pin `5559679229`); the SYNC path in the same file already loops all outputs (`serving_completion.cpp:286`), so the defect is specific to the async SSE source. Latent until [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s fan-out makes a second output reachable, so FIXED in that row's pull request. Spec: [async-parallel-sampling.md](../specs/async-parallel-sampling.md) | bug | -| [#2120](https://github.com/mudler/vllm.cpp/issues/2120) | — | **`/v1/chat/completions` streaming collapses `n > 1` onto one choice's parser and text state.** `ChatSseStream` (`src/vllm/entrypoints/openai/serving_chat.cpp:305-548`) holds `previous_text_`, `previous_num_tokens_` and `tools_streamed_` as scalars, ONE `parser_`/`engine_parser_`/`reasoning_parser_` instance for the whole response, emits the role frame for index 0 only (`:385`), and reads `response.outputs.front()`. Upstream keeps every one of those per choice index (`vllm/entrypoints/openai/chat_completion/serving.py:404-802` at pin `5559679229`). Repairing it is a parser-lifetime change with its own review surface, not a repair of the engine fan-out, so it is explicitly OUT of [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s scope. Owed under [async-parallel-sampling.md](../specs/async-parallel-sampling.md) `## Owed` | bug | -| [#2121](https://github.com/mudler/vllm.cpp/issues/2121) | — | **`AsyncLLM::add_request_wave` does not fan out `n > 1`.** Both overloads (`src/vllm/v1/engine/async_llm.cpp:133,166`) register every input with `request_index=0` and no `ParentRequest`, so a wave entry carrying `n > 1` is served as `n == 1` — the same defect [#1816](https://github.com/mudler/vllm.cpp/issues/1816) records for the single-request overloads. The wave is a LOCAL extension with no upstream counterpart and no OpenAI route reaches it; its only caller is `examples/bench/bench_core.h:222,225`, which is why it is excluded from #1816's fix rather than folded into it. Owed under [async-parallel-sampling.md](../specs/async-parallel-sampling.md) `## Owed` | bug | -| [#2145](https://github.com/mudler/vllm.cpp/issues/2145) | — | **The parallel-sampling fan-out DEEP-copies the prompt `n` times where upstream's copy is shallow.** Both fan-out sites build each child with `EngineCoreRequest child = request;` (`src/vllm/v1/engine/async_llm.cpp` `PublishParallelSampling`, `src/vllm/v1/engine/llm_engine.cpp` `FanOutParallelSampling`), and `EngineCoreRequest::prompt_token_ids` is a `std::vector` held BY VALUE (`include/vllm/v1/engine/types.h:79`), so each of the `n` children owns a full copy of the prompt and `Request::FromEngineCoreRequest` makes a second one per child — `O(n * prompt_len)` bytes moved before the first token is scheduled. Upstream copies ZERO prompt tokens: `copy(request)` (`vllm/v1/engine/async_llm.py:393`, `vllm/v1/engine/llm_engine.py:283` @ pin `5559679229`) is SHALLOW, every child references the same list object, and the last child reuses the parent outright. The comment on both of our lines claimed the copy "shares the prompt token ids", which is FALSE; [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s pull request corrects the comment and points here, and does NOT fix the cost, because the cheap mirror is a shared immutable token buffer on `EngineCoreRequest` that every engine path reads — a types-level change with its own review surface. No correctness effect; the cost scales with prompt length times `n`, so it is invisible on the short-prompt suites. Owed under [async-parallel-sampling.md](../specs/async-parallel-sampling.md) `## Owed` | bug | -| [#2150](https://github.com/mudler/vllm.cpp/issues/2150) | — | **`ParentRequest::get_outputs` indexes `output_aggregator_` unchecked and dereferences a possibly-empty optional; upstream raises `IndexError` where we get UB.** `src/vllm/v1/engine/parallel_sampling.cpp:75-82` writes `output_aggregator_[idx] = ...` and later drains with `*slot`, both 1:1 with `vllm/v1/engine/parallel_sampling.py:100-126` @ pin `5559679229`, which is equally unchecked — but upstream's aggregator is a Python list, so an out-of-range index raises and an unfilled slot surfaces as `None`. Ours is `std::vector>`, where both are undefined behaviour. NOT a live bug: `idx` is `0..n-1` by construction from `get_child_info` and the vector is sized `n`, so a bounds check today would be DEAD CODE and is deliberately not added. Recorded because this UB is why [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s `request_index=0` mutation has no stable exit status (RC=135/139 full, RC=1 case-scoped across rounds) and had to be recorded as a signal rather than a number; a debug-configuration assertion or `.at()` is the likely shape of a fix, not a release-path branch. Found in the fresh review of #1816. Owed under [async-parallel-sampling.md](../specs/async-parallel-sampling.md) `## Owed` | bug | +| [#2119](https://github.com/mudler/vllm.cpp/issues/2119) | `SAMPLE-N-ASYNC` | **`/v1/completions` streaming drops every choice past the first**: `CompletionSseStream::next` (`src/vllm/entrypoints/openai/serving_completion.cpp:92`) reads `response.outputs.front()` and formats one SSE choice from it, while `RequestOutputCollector::Merge` (`src/vllm/v1/engine/output_processor.cpp:96-118`) keeps distinct `index` completions side by side in ONE frame whenever the producer outruns the consumer. Upstream flattens the samples one choice per chunk and asserts it (`tests/entrypoints/openai/completion/test_completion.py:419-424`, `test_parallel_streaming` at pin `5559679229`); the SYNC path in the same file already loops all outputs (`serving_completion.cpp:286`), so the defect is specific to the async SSE source. Latent until [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s fan-out makes a second output reachable, so FIXED in that row's pull request. Spec: [async-parallel-sampling.md](specs/async-parallel-sampling.md) | bug | +| [#2120](https://github.com/mudler/vllm.cpp/issues/2120) | — | **`/v1/chat/completions` streaming collapses `n > 1` onto one choice's parser and text state.** `ChatSseStream` (`src/vllm/entrypoints/openai/serving_chat.cpp:305-548`) holds `previous_text_`, `previous_num_tokens_` and `tools_streamed_` as scalars, ONE `parser_`/`engine_parser_`/`reasoning_parser_` instance for the whole response, emits the role frame for index 0 only (`:385`), and reads `response.outputs.front()`. Upstream keeps every one of those per choice index (`vllm/entrypoints/openai/chat_completion/serving.py:404-802` at pin `5559679229`). Repairing it is a parser-lifetime change with its own review surface, not a repair of the engine fan-out, so it is explicitly OUT of [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s scope. Owed under [async-parallel-sampling.md](specs/async-parallel-sampling.md) `## Owed` | bug | +| [#2121](https://github.com/mudler/vllm.cpp/issues/2121) | — | **`AsyncLLM::add_request_wave` does not fan out `n > 1`.** Both overloads (`src/vllm/v1/engine/async_llm.cpp:133,166`) register every input with `request_index=0` and no `ParentRequest`, so a wave entry carrying `n > 1` is served as `n == 1` — the same defect [#1816](https://github.com/mudler/vllm.cpp/issues/1816) records for the single-request overloads. The wave is a LOCAL extension with no upstream counterpart and no OpenAI route reaches it; its only caller is `examples/bench/bench_core.h:222,225`, which is why it is excluded from #1816's fix rather than folded into it. Owed under [async-parallel-sampling.md](specs/async-parallel-sampling.md) `## Owed` | bug | +| [#2145](https://github.com/mudler/vllm.cpp/issues/2145) | — | **The parallel-sampling fan-out DEEP-copies the prompt `n` times where upstream's copy is shallow.** Both fan-out sites build each child with `EngineCoreRequest child = request;` (`src/vllm/v1/engine/async_llm.cpp` `PublishParallelSampling`, `src/vllm/v1/engine/llm_engine.cpp` `FanOutParallelSampling`), and `EngineCoreRequest::prompt_token_ids` is a `std::vector` held BY VALUE (`include/vllm/v1/engine/types.h:79`), so each of the `n` children owns a full copy of the prompt and `Request::FromEngineCoreRequest` makes a second one per child — `O(n * prompt_len)` bytes moved before the first token is scheduled. Upstream copies ZERO prompt tokens: `copy(request)` (`vllm/v1/engine/async_llm.py:393`, `vllm/v1/engine/llm_engine.py:283` @ pin `5559679229`) is SHALLOW, every child references the same list object, and the last child reuses the parent outright. The comment on both of our lines claimed the copy "shares the prompt token ids", which is FALSE; [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s pull request corrects the comment and points here, and does NOT fix the cost, because the cheap mirror is a shared immutable token buffer on `EngineCoreRequest` that every engine path reads — a types-level change with its own review surface. No correctness effect; the cost scales with prompt length times `n`, so it is invisible on the short-prompt suites. Owed under [async-parallel-sampling.md](specs/async-parallel-sampling.md) `## Owed` | bug | +| [#2150](https://github.com/mudler/vllm.cpp/issues/2150) | — | **`ParentRequest::get_outputs` indexes `output_aggregator_` unchecked and dereferences a possibly-empty optional; upstream raises `IndexError` where we get UB.** `src/vllm/v1/engine/parallel_sampling.cpp:75-82` writes `output_aggregator_[idx] = ...` and later drains with `*slot`, both 1:1 with `vllm/v1/engine/parallel_sampling.py:100-126` @ pin `5559679229`, which is equally unchecked — but upstream's aggregator is a Python list, so an out-of-range index raises and an unfilled slot surfaces as `None`. Ours is `std::vector>`, where both are undefined behaviour. NOT a live bug: `idx` is `0..n-1` by construction from `get_child_info` and the vector is sized `n`, so a bounds check today would be DEAD CODE and is deliberately not added. Recorded because this UB is why [#1816](https://github.com/mudler/vllm.cpp/issues/1816)'s `request_index=0` mutation has no stable exit status (RC=135/139 full, RC=1 case-scoped across rounds) and had to be recorded as a signal rather than a number; a debug-configuration assertion or `.at()` is the likely shape of a fix, not a release-path branch. Found in the fresh review of #1816. Owed under [async-parallel-sampling.md](specs/async-parallel-sampling.md) `## Owed` | bug | | [#1756](https://github.com/mudler/vllm.cpp/issues/1756) | `ENG-PUBLIC-DOC-PROJECTIONS` | Three user-facing benchmark references retained the old public-document shape after #1714: README sent two claims to the internal benchmark record, the speculative-decoding guide described the compact index as the old record, and the llama.cpp detail page named deleted `docs/STATUS.md`. Fixed in flow by linking the public benchmark detail pages directly. No result, support claim, or README news headline changed | bug | -| [#2151](https://github.com/mudler/vllm.cpp/issues/2151) | `SPEC-DFLASH2` | **`cb28167c9` (#2148) reverted the reorder-threshold wiring on an uncontrolled measurement, so the revert comes back out.** The revert's stated reason — the wiring "costs 38% at c=8" — compared two builds run hours apart, in commit order, on a drifting box, with neither arm re-measured. An interleaved A/B with a terminal control (`A B A B A B A`, one lease, one hour, 2026-08-28) settles it: the arm WITHOUT the wiring (`16ebcac4b`) read 56.22, 51.29 and 36.82 out tok/s, and the arm WITH it (`5e9d81dad`) read 34.66, 35.49, 43.30 and 78.86 — the wired arm holds both the LOWEST and the HIGHEST reading in the set, so no ordering between the builds exists and the 38% was an artifact of the instrument. The wiring is justified without any throughput claim: `SpecAsDecodeReorderThreshold` (`include/vllm/v1/attention/backend.h:180-184`) mirrors upstream's `1 + (parallel_drafting ? 2 : 1) * k` and reaches only the spec-as-decode classification at `:201`, while the reorder the value exists to bound takes the declaration default of 1, against upstream's `_may_reorder_batch` passing `decode_threshold=self.reorder_batch_threshold` (`gpu_model_runner.py:1126-1130` @ pin `5559679229`). Restores #2138 byte-for-byte and repairs `## Now`, `## Outcome` and `## Owed` in [specs/reorder-threshold-wiring.md](../specs/reorder-threshold-wiring.md), whose `## WITHDRAWN` section carried the same false premise. The instrument defect the run exposed is larger than this row and is tracked separately by [#2152](https://github.com/mudler/vllm.cpp/issues/2152) | bug | -| [#2152](https://github.com/mudler/vllm.cpp/issues/2152) | — | **The c=8 ladder rung has a 127% spread, so every single-run comparison at that rung is ungated, including the standing vLLM and SGLang positions.** Seven interleaved runs on 2026-08-28, two builds, one lease, one hour, every arm re-measured: `16ebcac4b` read 56.22 / 51.29 / 36.82 out tok/s and `5e9d81dad` read 34.66 / 35.49 / 43.30 / 78.86. One UNCHANGED binary spans 52%, the other 127%. The instrument's spread is larger than every effect it has been asked to detect. The 5.9% c=8 figure quoted throughout this repository comes from a 4-run study that sampled a stable window and has since been used as though it bounded the rung; it does not, and a number quoted often became treated as measured. Everything gated at c=8 with n=1 per arm is therefore ungated: #2148's 38% (void, see [#2151](https://github.com/mudler/vllm.cpp/issues/2151)), the W12 and W13 c=8 attributions, and the "parity with vLLM, 23% behind SGLang" position. Owes three things — a repeat count DERIVED from the measured spread rather than assumed, interleaved arms plus a terminal control in the harness itself so a drifting box invalidates its own run instead of returning a confident number, and a cause for the drift (a clock pin outliving a lease is the first hypothesis and is untested). Owed under `## Owed` in [specs/reorder-threshold-wiring.md](../specs/reorder-threshold-wiring.md) until a row picks it up | bug | -| [#2107](https://github.com/mudler/vllm.cpp/issues/2107) | `BACKEND-TENSTORRENT-QWEN35` | **TT eager decode spends the wall on host staging, not device kernels.** The #1715 profile lever measured one Qwen3.5-0.8B eager step at 0.104 tok/s and ranked NO device kernel: `vt::Tensor::Numel()` 27.09% of samples, `EnsureDevice2D`→`MatmulBTKernel` 24.10% of the call graph feeding `DenseMlpBlock` 9.44% and the f32 logits GEMM 9.32%, repeated TT-Metal context/UMD discovery ~12%, threadpool spin 11.4%, `memcpy` 7.04%, `bfloat16::from_float` 2.62%; the same leg without the TT backend runs 7.521 tok/s (~73x). Root cause is structural, not a missing kernel: `EnsureDevice2D` (`src/vt/tenstorrent/tenstorrent_ops.cpp:434`) stages element-by-element through an f32 intermediate (`LoadElemF32` per index, then f32→bf16 again in `UploadRows`), takes up to four `FindSlot` mutex acquisitions per call, and every upload re-resolves TT-Metal context/device/chip handles. The three levers are the recorded next gate on the open-gaps row: cache resolved handles, hoist shape math and bulk the element loop, batch per-layer staging. Constraint: numerics may not move — sacred golden 16/16 and the full TT suite stay green; captured tracing stays blocked behind #1625. Evidence: `docs/bench-evidence/tt-qwen35-eager-profile-20260827.log`, [record](../benchmark-record.md) | perf | +| [#2151](https://github.com/mudler/vllm.cpp/issues/2151) | `SPEC-DFLASH2` | **`cb28167c9` (#2148) reverted the reorder-threshold wiring on an uncontrolled measurement, so the revert comes back out.** The revert's stated reason — the wiring "costs 38% at c=8" — compared two builds run hours apart, in commit order, on a drifting box, with neither arm re-measured. An interleaved A/B with a terminal control (`A B A B A B A`, one lease, one hour, 2026-08-28) settles it: the arm WITHOUT the wiring (`16ebcac4b`) read 56.22, 51.29 and 36.82 out tok/s, and the arm WITH it (`5e9d81dad`) read 34.66, 35.49, 43.30 and 78.86 — the wired arm holds both the LOWEST and the HIGHEST reading in the set, so no ordering between the builds exists and the 38% was an artifact of the instrument. The wiring is justified without any throughput claim: `SpecAsDecodeReorderThreshold` (`include/vllm/v1/attention/backend.h:180-184`) mirrors upstream's `1 + (parallel_drafting ? 2 : 1) * k` and reaches only the spec-as-decode classification at `:201`, while the reorder the value exists to bound takes the declaration default of 1, against upstream's `_may_reorder_batch` passing `decode_threshold=self.reorder_batch_threshold` (`gpu_model_runner.py:1126-1130` @ pin `5559679229`). Restores #2138 byte-for-byte and repairs `## Now`, `## Outcome` and `## Owed` in [specs/reorder-threshold-wiring.md](specs/reorder-threshold-wiring.md), whose `## WITHDRAWN` section carried the same false premise. The instrument defect the run exposed is larger than this row and is tracked separately by [#2152](https://github.com/mudler/vllm.cpp/issues/2152) | bug | +| [#2152](https://github.com/mudler/vllm.cpp/issues/2152) | — | **The c=8 ladder rung has a 127% spread, so every single-run comparison at that rung is ungated, including the standing vLLM and SGLang positions.** Seven interleaved runs on 2026-08-28, two builds, one lease, one hour, every arm re-measured: `16ebcac4b` read 56.22 / 51.29 / 36.82 out tok/s and `5e9d81dad` read 34.66 / 35.49 / 43.30 / 78.86. One UNCHANGED binary spans 52%, the other 127%. The instrument's spread is larger than every effect it has been asked to detect. The 5.9% c=8 figure quoted throughout this repository comes from a 4-run study that sampled a stable window and has since been used as though it bounded the rung; it does not, and a number quoted often became treated as measured. Everything gated at c=8 with n=1 per arm is therefore ungated: #2148's 38% (void, see [#2151](https://github.com/mudler/vllm.cpp/issues/2151)), the W12 and W13 c=8 attributions, and the "parity with vLLM, 23% behind SGLang" position. Owes three things — a repeat count DERIVED from the measured spread rather than assumed, interleaved arms plus a terminal control in the harness itself so a drifting box invalidates its own run instead of returning a confident number, and a cause for the drift (a clock pin outliving a lease is the first hypothesis and is untested). Owed under `## Owed` in [specs/reorder-threshold-wiring.md](specs/reorder-threshold-wiring.md) until a row picks it up | bug | +| [#2107](https://github.com/mudler/vllm.cpp/issues/2107) | `BACKEND-TENSTORRENT-QWEN35` | **TT eager decode spends the wall on host staging, not device kernels.** The #1715 profile lever measured one Qwen3.5-0.8B eager step at 0.104 tok/s and ranked NO device kernel: `vt::Tensor::Numel()` 27.09% of samples, `EnsureDevice2D`→`MatmulBTKernel` 24.10% of the call graph feeding `DenseMlpBlock` 9.44% and the f32 logits GEMM 9.32%, repeated TT-Metal context/UMD discovery ~12%, threadpool spin 11.4%, `memcpy` 7.04%, `bfloat16::from_float` 2.62%; the same leg without the TT backend runs 7.521 tok/s (~73x). Root cause is structural, not a missing kernel: `EnsureDevice2D` (`src/vt/tenstorrent/tenstorrent_ops.cpp:434`) stages element-by-element through an f32 intermediate (`LoadElemF32` per index, then f32→bf16 again in `UploadRows`), takes up to four `FindSlot` mutex acquisitions per call, and every upload re-resolves TT-Metal context/device/chip handles. The three levers are the recorded next gate on the open-gaps row: cache resolved handles, hoist shape math and bulk the element loop, batch per-layer staging. Constraint: numerics may not move — sacred golden 16/16 and the full TT suite stay green; captured tracing stays blocked behind #1625. Evidence: `docs/bench-evidence/tt-qwen35-eager-profile-20260827.log`, [record](benchmark-record.md) | perf | | [#2115](https://github.com/mudler/vllm.cpp/issues/2115) | `BACKEND-TENSTORRENT-QWEN35` | **The `VT_TT_HOST_FREE_DECODE=0` leg of the Qwen3.5-0.8B sacred e2e drifts one anchor token — deterministic and pre-existing on main.** prompt[2] tok=1: engine 15039 vs committed anchor 1814, byte-identical across base `0ac84a486` (W4 stashed, rebuilt) and the W4 worktree twice. The ambient leg is 16/16 PASS / 0 forward-divergent / max gap 375 mnats; only the opt-out arm drifts. Found by the W4 gate run (#2107) and proven not a W4 regression; something between the `c31cad9c1` golden re-derivation and `8f5d4e4ed` moved the opt-out arm (candidates: `2a42cb369` KV-GDN-STATE-BUDGET, the W2c residency chain — unbisected). Owed: bisect to the landing that moved it, then fix the engine defect or re-derive the golden through the sanctioned procedure with the justification recorded — never to silence the gate. Evidence: `docs/bench-evidence/tt-qwen35-eager-leg2-anchor-drift-20260827.log` | bug | | [#2156](https://github.com/mudler/vllm.cpp/issues/2156) | `MODEL-MM-QWEN4-EXP` | **The PLE dilated depthwise causal conv had no `vt::` op, and neither half of what it needs exists on one.** Measured at `d598ffd77` rather than asserted: `CausalConv1dFwd` is causal and carries a persistent `conv_state`, and `CausalConv1dArgs` has NO dilation field at all (only `silu_activation`, `batch_ptr`, `token_chunk_offset_ptr`); `DepthwiseConv1d` / `Conv1d` / `ConvTranspose1d` dilate and are stateless. PLE needs both — `kernel_size = 4`, `dilation = ngram_size = 3`, so output `t` reads lags {9, 6, 3, 0} and the state is a genuine 9-deep history read at stride 3, `(K-1)*dilation` and not `K-1`. **Closed by a NEW OpId rather than by widening `CausalConv1dArgs`, on evidence:** that struct is read by five backends (`cpu_ops.cpp`, `cuda_gdn.cu`, `rocm_gdn_conv.hip`, `vulkan_ops.cpp`, `tenstorrent_ops.cpp`) across three OpIds, a new field is silently IGNORED by every kernel that does not read it, and four of those five arms cannot be gated on a CPU-only host — so the field would ship as a live wrong-answer path on the Mamba/GDN/KDA/Kimi conv instead of a refusal; and the `K-1` state width is welded into the spec-decode path's `max_query_len = state_len - (k - 1) + 1` (`src/vt/ops.cpp:1997`), which turns a nine-column dilated state at K = 4 into a meaningless bound of 7 and feeds it to the per-request checks at `:2007` and `:2009`, so widening it would weaken the guard for every existing caller. (The shared `CheckConvCommon` is NOT that weld and this row does not claim it is: its width test at `src/vt/ops.cpp:1732` is `conv_state.shape[2] >= k - 1`, a lower bound already widened for spec-decode, which a nine-column state passes at 9 >= 3.) Proven non-invasive: `CausalConv1dArgs` is untouched and `test_ops_gdn` (62), `test_gdn_prefill_conv` (10), `test_conv_update_fast` (1) and `test_ops_mamba2_state_update` (6) stay green. `vt::Qwen4ExpPleConv` batches over sequences through `query_start_loc`, takes a nullable per-sequence cache-row index, and has NO `has_initial_state`, because upstream's first call left-zero-pads (`cache_utils.py:1053-1060`) and a zeroed row is bit-identical to that. CPU kernel in `src/vt/cpu/cpu_qwen4_exp_ple.cpp`, gated against the SAME lane-pinned transformers v5.16.0 goldens the W2 host reference answers to, plus a 10240-channel model-width `memcmp` agreement check. **The dilation is gated as a VARIABLE, not assumed**: the generator runs upstream's own `_short_conv` at dilations 3, 2 and 1 over the same input and the same weight, the three answers separate by 0.443 to 0.716 against a 1e-5 tolerance, and the test re-measures that separation so a future regeneration that collapsed it could not pass in silence. 16-mutation battery, 15 RED, each proved applied by sha256 and each build rc read before any test result — five mutations first failed to BUILD under `-Werror` and would have run a stale binary as a pass. The one survivor, M10, is an EQUIVALENT MUTANT with the equivalence proved in the spec, and it is repaired by M16 rather than recorded. **UNREACHED at this merge commit and deliberately so**, exactly as W2/W3/W4/W5b-2 landed: the only production entry point is `ModelRegistry::Forward`, which is all-or-nothing; the wiring is owed by W5b under [#2031](https://github.com/mudler/vllm.cpp/issues/2031) and tracked by [#1978](https://github.com/mudler/vllm.cpp/issues/1978), and reaching this op from the runner's recurrent cache additionally waits on [#2131](https://github.com/mudler/vllm.cpp/issues/2131). Also owed and named in the spec's `## Owed`: the CUDA arm and the accumulator width it must choose, and a bf16 `conv_state` the dispatcher refuses until an arm can write one | feature | | [#2158](https://github.com/mudler/vllm.cpp/issues/2158) | `SERVE-CHAT-TEMPLATE` | The GGUF chat-template selection that #2079 wires is gated as a FUNCTION and not as the server path: deleting the `LoadChatTemplateForModel` call site in `server_main.cpp` leaves `test_chat_template` at 37 cases / 147 assertions green, so the #2077 regression could return unseen. MEASURED on the #2079 head merged onto main, not argued. Landed with the gap named because the wiring is five lines at a production entry point and was verified by hand on gfx1100, while the defect makes every GGUF chat request useless. Owed: a case entering through `VllmServerMain` in the `test_serve_residency_config.cpp` re-exec shape. Its obstacle is why this is its own unit of work — the chat-template block sits after the full engine load, so the nonexistent-model-directory trick cannot reach it and the synthetic GGUF stops one step earlier at the missing tokenizer | bug | @@ -841,73 +669,63 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#2166](https://github.com/mudler/vllm.cpp/issues/2166) | `ENG-MM-INPUT-PIPELINE` | **Muse Glimmer's tower could not ride the #1359 bf16 storage change, because its `compute_dtype = kF32` per-stage gate computes on the stored weight values.** The spec's §4.4 ruling that widening back is bit-identical holds for the production loader reading an all-BF16 checkpoint; it does not hold for `test_muse_glimmer_vision`, whose weights are a synthetic f32 LCG that `scripts/mm/muse_glimmer_vision_ref.py:52-61` builds as `torch.float32` and never rounds. MEASURED on a scratch tree with the bf16 store applied: the five f32-arm stages move from rel_l2 1.0-3.0e-07 to 2.164e-03 / 2.193e-03 / 2.220e-03 / 2.892e-03 / 3.462e-03 against a 1e-6 bound — five assertions red, three orders out. The PRODUCTION path is unaffected and that is measured too: the bf16 arm read `rel_l2=5.951e-03 max_abs=3.675e-02` in the same tree, byte-for-byte what it reads today, so the 3.580 GiB -> 7.161 GiB widening is genuinely removable and only the gate stands in the way. Owed: round the LCG through bf16 on BOTH sides and regenerate `muse_glimmer_vision_goldens.inc` (a reference change that needs its own red/green argument and must not ride in the change it gates), then narrow the four structs, grow `Upload` into the `UploadWeight` shape Qwen3-VL now has, and restore `TOWER_RESIDENT_BYTES` for `muse-glimmer` plus the `WIDEN` mirror to 1x. Threshold unchanged from `specs/vision-tower-dtype-polarity.md` §6.1: >= 3,459,322,368 B on the default arm | bug | | [#2164](https://github.com/mudler/vllm.cpp/issues/2164) | `BACKEND-ROCM` | **gfx1100 decode is launch-bound, and kernel micro-optimization is exhausted as a lever.** The GFX1100-TG200 campaign swept 15 levers, adopted 11, and reached ~103 tok/s (~9.71 ms/tok) against a 200 tok/s target on Qwen3.5-4B-Q4_K_M / RX 7900 XTX. The discriminating result is a NEGATIVE one: T20's full-warp `KQuantGemvMmvqRow` rewrite is 2.38x-3.13x faster on large grids in microbenchmark, and a paired interleaved 5-rep engine A/B reads 92.9 vs 92.8 tok/s — a 0.1% wash — because the dominant Q4_K path runs at grid ~576 and is bound by fixed launch cost, not by the reduction barriers the rewrite removed, while the large-grid win lands on lm_head at one call per token (~0.04 ms/tok averaged). An earlier `rocprofv3` capture shows the mechanism directly: 97 standalone `QuantizeQ8KK` launches per token, EVERY one a single block (`m*nsb <= 128` at batch 1), mean duration 48.2-50.1 us FLAT with respect to K. The issue also separates two overhead terms the campaign's summary collapses: ~4.2 ms/tok is kernel time above the 4.38 ms/tok weight-read floor (occupancy and per-launch cost INSIDE kernels), and a further ~1.13 ms/tok is wall outside kernels entirely. Next levers are HIP graph capture ([#332](https://github.com/mudler/vllm.cpp/issues/332), which predicted this on gfx1200 from an explicitly unmeasured two-point fit and which this measures on gfx1100), a `SiluMulK` quant epilogue for the 40 of 97 launches the `RmsNorm` epilogue cannot absorb, then persistent kernels. Owed: the evidence is read from unmerged [#1936](https://github.com/mudler/vllm.cpp/pull/1936) at `b058bb752` and is NOT reproducible from `main`, so landing `docs/bench-evidence/gfx1100-tg200-*.md` and the campaign spec comes first, then a fresh capture with per-token dispatch counts to replace the budget-table arithmetic with a traced split | perf | | [#2167](https://github.com/mudler/vllm.cpp/issues/2167) | `MODEL-MM-QWEN4-EXP` | **Qwen Sparse Attention on the device arm, and the finding that HALF of it was already a `vt::` op.** The issue's "why nothing existing serves it" table named `IndexSelect`, `TopKValuesIndices`, `GatherMlaCache` and the fused `kDeepseekV4Dsa`/`kDeepseekV4Compressor`, and it omitted the two that DO serve: `vt::DsaIndexerLogits` computes `sum_h fold[t,h] * ReLU(dot(q[t,h,:], k[s,:]))` over a one-key-head MQA cache with a per-query `[win_start, win_end)` window, and with `weights` all ones, a null `q_scale` and `n_head_scale = 1` its fold collapses to the single constant `softmax_scale` — set that to `index_head_dim ** -0.5` and it IS `Qwen4ExpTextQSAIndexer`'s block score, QSA having neither DeepSeek-V4's learned `weights_proj` nor its `n_head ** -0.5`; and `vt::DsaTopkSelect` is the same all-select-below-k, ties-to-the-LOWER-index, ASCENDING-emission top-k over the block axis. So the indexer is COMPOSED from those two rather than re-implemented beside `cpu_dsa_indexer.cpp`, which would have been the parallel path AGENTS.md §"Shared seams" forbids. Measured, not argued: the composed selection equals the transformers v5.16.0 selected-token sets for every query token of both fixtures, ragged tail included. **Two ops are genuinely new**: `vt::Qwen4ExpQsaCompress` (mean pool over a NON-overlapping window of `compress_ratio`, `k_layernorm` on the pooled key, RoPE at the block's FIRST token) — the pool has no `vt::` counterpart at all, this tree having no mean, no pool, no axis reduction and no transpose to fake one with, and fusing the three stages mirrors upstream's own single kernel on the `kFusedNormRope` precedent; and `vt::Qwen4ExpQsaGatherAttention`, which has no counterpart anywhere, every DeepSeek-V4 sparse consumer attending the COMPRESSED MLA KV and MiniMax-M3's attending raw tokens at KV-PAGE granularity while QSA attends RAW tokens at ratio-4 granularity. CPU kernels in `src/vt/cpu/cpu_qwen4_exp_qsa.cpp`. **The gates run PAST the 2048-token indexer budget**, which the spec's `## Gates` requires because below it every candidate is selected: the released indexer config at 3002 tokens of context discards 238 of 750 complete blocks, and the 2051-token control beside it selects everything and is the measurement of why the requirement exists. **A mutation survived and was REPAIRED rather than recorded**, and it is the one this wave exists to stop: a dense masked walk over every cached row that reports the SPARSE `keys_visited` figure passed 10 cases / 4167 assertions, exactly as W4's M22c passed 12 / 7251 — a counter a kernel writes cannot convict the kernel that writes it, and no value comparison can either, because `exp(-inf - m)` is exactly +0 so a mask agrees with a gather value for value. The repair is an observable of the WALK: a cache whose UNSELECTED rows are NaN, which a gather never addresses and a mask multiplies by a zero weight into `0.0f * NaN` = NaN. **UNREACHED at this merge commit and deliberately so**, as W2/W3/W4/W5b-1/W5b-2 landed: `ModelRegistry::Forward` is the only production entry point and is all-or-nothing; the wiring is owed by W5b under [#2031](https://github.com/mudler/vllm.cpp/issues/2031), tracked by [#1978](https://github.com/mudler/vllm.cpp/issues/1978), and reaching the ops from the runner's caches additionally waits on [#2131](https://github.com/mudler/vllm.cpp/issues/2131). Also owed and named in the spec's `## Owed`: the CUDA arm of both ops and the three decisions it must make first, a single-pass online softmax, the ungated bf16 STORAGE arms, and the paged side-cache store | feature | -| [#2173](https://github.com/mudler/vllm.cpp/issues/2173) | `MODEL-MM-gemma4-mm-gemma4-for-conditional-generation` | **The Gemma-4 SigLIP2 vision tower has no production caller, so every measurement quoted about it measures a class rather than a capability.** Measured at `a1dcc74f4`: `grep -rn 'gemma4_vision.h' src/ include/` returns exactly one hit outside the header, `gemma4_vision.cpp:16` including its own header, and the only other includers are `tests/vllm/multimodal/test_gemma4_vision_tower.cpp:25` and `test_gemma4_registry_e2e.cpp:42`; every `Gemma4VisionForward` / `Gemma4VisionWeights` call site is inside those two files. The tower is unreached TWICE OVER, which is what distinguishes it from #1358: the engine driver `Gemma4GenerateGreedyViaRegistry` (`gemma4_mm.cpp:165`) takes `mm_projected` as a CALLER-SUPPLIED `const std::vector&` and masked-scatters it at `:250-252` without ever calling the tower, and that driver's own only caller is `test_gemma4_registry_e2e.cpp:244`. Per [`.agents/reachability.md`](../reachability.md) this is the test-only-driver shape, and a change with no production call site to delete has already answered the question — there is nothing to mutate. The 2026-07-29 `MM-IMAGE-E2E` fold recorded on the owning row IS real at the `ModelRegistry::Forward` layer (`gemma4_registry.cpp:151` routes `ModelForwardInput.mm` into `Gemma4Model::ForwardMm`); what is missing is everything above it that would build an `mm` field for Gemma-4 from an image. Consequence already observed: #2169's body claimed "Gemma-4 ran that pass on every image" about a per-weight `F32ToBF16` upload pass that no image reaches, and the operator repeated it. Third instance of a class whose other two are filed — #1358 (Qwen3-VL loads its tower and never reads it back) and #1566 (Muse Glimmer's encoder has no production caller) — and the one that was undisclosed. Filed from the fresh review of #2169; listed under `## Owed` in [`vision-tower-dtype-polarity.md`](../specs/vision-tower-dtype-polarity.md) | bug | +| [#2173](https://github.com/mudler/vllm.cpp/issues/2173) | `MODEL-MM-gemma4-mm-gemma4-for-conditional-generation` | **The Gemma-4 SigLIP2 vision tower has no production caller, so every measurement quoted about it measures a class rather than a capability.** Measured at `a1dcc74f4`: `grep -rn 'gemma4_vision.h' src/ include/` returns exactly one hit outside the header, `gemma4_vision.cpp:16` including its own header, and the only other includers are `tests/vllm/multimodal/test_gemma4_vision_tower.cpp:25` and `test_gemma4_registry_e2e.cpp:42`; every `Gemma4VisionForward` / `Gemma4VisionWeights` call site is inside those two files. The tower is unreached TWICE OVER, which is what distinguishes it from #1358: the engine driver `Gemma4GenerateGreedyViaRegistry` (`gemma4_mm.cpp:165`) takes `mm_projected` as a CALLER-SUPPLIED `const std::vector&` and masked-scatters it at `:250-252` without ever calling the tower, and that driver's own only caller is `test_gemma4_registry_e2e.cpp:244`. Per [`.agents/reachability.md`](reachability.md) this is the test-only-driver shape, and a change with no production call site to delete has already answered the question — there is nothing to mutate. The 2026-07-29 `MM-IMAGE-E2E` fold recorded on the owning row IS real at the `ModelRegistry::Forward` layer (`gemma4_registry.cpp:151` routes `ModelForwardInput.mm` into `Gemma4Model::ForwardMm`); what is missing is everything above it that would build an `mm` field for Gemma-4 from an image. Consequence already observed: #2169's body claimed "Gemma-4 ran that pass on every image" about a per-weight `F32ToBF16` upload pass that no image reaches, and the operator repeated it. Third instance of a class whose other two are filed — #1358 (Qwen3-VL loads its tower and never reads it back) and #1566 (Muse Glimmer's encoder has no production caller) — and the one that was undisclosed. Filed from the fresh review of #2169; listed under `## Owed` in [`vision-tower-dtype-polarity.md`](specs/vision-tower-dtype-polarity.md) | bug | | [#2174](https://github.com/mudler/vllm.cpp/issues/2174) | `ENG-MM-INPUT-PIPELINE` | **`MakeDevBf16` in `gemma4_vision.cpp` allocated from the tensor's declared SHAPE and copied the host store's OWN length into it, with nothing checking that the two agree.** `bytes = numel * SizeOf(kBF16)` sizes the allocation while `b.Copy(q, d.p, bf.data(), bf.size() * sizeof(uint16_t))` sets the copy length, so `bf.size() * 2 > bytes` overruns the allocation and a loader bug — a wrong enumeration, a mis-shaped weight, a checkpoint whose config disagrees with its tensors — lands as heap corruption rather than as a named refusal; the under-full case leaves an uninitialised tail. Its line-for-line twin `qwen3_vl_vision.cpp:137` grew exactly this guard in #1359 and copies `bytes`, so the asymmetry is the defect and the Gemma-4 copy predates #1359 rather than being made worse by it. Not currently exploitable for a reason that is itself debt: nothing in production calls `Gemma4VisionForward` (#2173), so the only shapes this function sees are two tests' fixtures, which agree by construction. Found in the fresh review of #2169 and FIXED IN FLOW there by mirroring the twin's `VT_CHECK` and copying `bytes`; behaviour is unchanged on every shape the loaders produce | bug | | [#2171](https://github.com/mudler/vllm.cpp/issues/2171) | `SPEC-DFLASH2` | **`DFlashAttnMmaKernel`'s multi-block QUERY path has never executed under test: every D1 case stops at `Tq=27` and the tile is 64 rows.** The kernel tiles the query axis at `kMmaWarps * kMmaQ = 4 * 16 = 64` (`src/vt/cuda/cuda_ops.cu:2372-2373`, grid `:2676`), and `RunD1Bf16Parity`'s four cases in `tests/vt/test_ops_dflash_block_attn.cpp` carry `Tq` of 18, 18, 18 and 27, so `mgrid.x` has always been 1. Production crosses the boundary on EVERY step — 8 concurrent requests at k=8 is `Tq = 8*9 = 72`, two query blocks, the second holding only the last request's nine rows. The comment above those cases reasons about walking several `kMmaKeys` tiles, which is the KEY axis; the query axis had no coverage past its first block. THIRD instance of this shape in one file, one axis over each time: a tiled CUDA path guarded to `num_reqs == 1` that "shipped never-executed while the suite stayed green", then an f32 harness that "by dispatch can never reach `DFlashAttnMmaKernel`" (the reason `RunD1Bf16Parity` exists). **The kernel PASSES at the missing shapes** — six added cases run on dgx:gpu0 GB10 sm_121a give 10 cases / 89886 assertions / ZERO failures, max\|diff\| 1.3e-4 — so this is a coverage gap, not a live defect, and a future regression there would have landed green. Controls are chosen for ATTRIBUTION: an 8-request red beside a 7-request `Tq=63` green isolates the query-block boundary, and a production-scale red (8 reqs, ctx ~1200, `Ncomb` ~9.7k) beside a single-request control at the same key extent isolates the many-request key union from context length. Found while investigating [#2154](https://github.com/mudler/vllm.cpp/issues/2154), where the query tile was a candidate mechanism for the acceptance collapse; these cases REFUTED that hypothesis | bug | | [#2176](https://github.com/mudler/vllm.cpp/issues/2176) | `MODEL-MM-dots3-note` | **dots3-note's nextn refusal is STRICTER than vLLM, which DROPS `model.layers.46.*` and `model.mtp.*` from the main model rather than refusing.** `Dots3NoteDeviceRefusal` turns away any config with `num_nextn_predict_layers > 0`, and §4 trap 3 correctly defaults that to 1 for a released `config.json` that does not carry the key — so every released checkpoint trips a branch upstream does not have. vLLM skips those weights in three places, re-derived at the row's pin `bc2d63e650`: `utils.py:542` `get_spec_layer_idx_from_weight_name` (matching `model.layers.{base+i}.` at `:559`), `deepseek_v2.py:1618-1620` `if spec_layer is not None: continue # skip spec decode layers for main model`, and `models/dots3_note/nvidia/model.py:624` `if name.startswith("mtp."): continue` inside `Dots3NoteModel._adapt_weights`. `Dots3NoteLanguageModelForCausalLM` (`model.py:681`) subclasses `DeepseekV32ForCausalLM`, so the second is the path this architecture loads through. The repair is the classifier-deferral shape the vision and audio towers already use: a `nextn` bucket on `Dots3NoteAccounting`, filled by a `Dots3NoteIsNextnTensor(params, name)` predicate rather than a static-prefix table row because the prefix is config-derived, with the 19 tensors staying ENUMERATED so an absent one still refuses. Over the released index the split becomes 35362 language / 19 nextn / 2195 vision / 430 audio = 38006, against W2's 35381 / 2195 / 430. Fixed in flow with W5 (the MoE brick), because W5 is what makes the other half of the released config representable and the two together are what let `Dots3NoteDeviceRefusal(released_params)` return empty for the first time. Not the MTP head, which stays W10 | bug | -| [#2190](https://github.com/mudler/vllm.cpp/issues/2190) | `MODEL-MM-dots3-note` | **dots3-note's quantization refusal keys on `weight_block_size` alone, so a per-tensor or per-row fp8/gptq/awq config loads and silently dequantizes to bf16.** `Dots3NoteDeviceRefusal` (`dots3_note_device.cpp:855`) branches on `has_blockwise_quant()`, which is `!weight_block_size.empty()` (`dots3_note.h:208`). `quant_method` IS parsed (`dots3_note.cpp:264-268`) and stored (`dots3_note.h:206`), and is read for nothing but the text of the blockwise message (`:862-863`). A `config.json` with `quantization_config.quant_method = "fp8"` (or gptq/awq/mxfp4/compressed-tensors) and NO `weight_block_size` therefore passes, and `dense_loaders::MaterializeBf16Source` silently dequantizes a per-tensor or per-output-ROW `_scale` into a bf16 GEMM — which is precisely the case the refusal's own comment names as the worse one, five lines above the branch that does not cover it (`:849-854`). This row has NO oracle on any hardware we own (spec §6.4), so nothing downstream catches the plausible wrong answer. No released checkpoint is affected: the bf16 repo carries no `quantization_config` and the `-fp8` sibling carries `weight_block_size [128, 128]` and is refused correctly. Owed: refuse a non-empty `quant_method` this port cannot read, naming the method and W9, with the config-fixture gate the blockwise case already has. Found by the fresh review of [#2187](https://github.com/mudler/vllm.cpp/pull/2187) as F5 and deliberately not fixed there — a refusal-semantics change needs its own red-before fixture. Under `## Owed` in [specs/dots3-note.md](../specs/dots3-note.md) | bug | -| [#2191](https://github.com/mudler/vllm.cpp/issues/2191) | `MODEL-MM-dots3-note` | **`hidden_act` is not mirrored, so a non-silu dots3-note config runs SwiGLU silently where vLLM raises.** `DeepseekV2MoE.__init__` refuses anything but silu before it builds a thing — `deepseek_v2.py:310-314` @ pin `bc2d63e650`, `ValueError(f"Unsupported activation: {config.hidden_act}. Only silu is supported for now.")` — and `Dots3NoteLanguageModelForCausalLM` (`model.py:681`) subclasses `DeepseekV32ForCausalLM`, so that is the `__init__` W5 ports. `grep -c hidden_act` over `dots3_note.cpp` and `dots3_note.h` is **0**: the key is never parsed, `Dots3NoteParams` has no field for it, and `Dots3NoteDeviceRefusal` never mentions it, so `hidden_act: "gelu"` loads and runs `vt::MoeSiluMul` / `vt::MoeGroupedGemmBf16GateUpSilu` with no refusal. Same hole in `deepseek_v2.cpp` and `deepseek_v2_weights.cpp` (`grep -c` = 0 on both), so it is a MIRROR GAP inherited by both ports rather than a W5 regression; `parakeet_transducer.cpp:108-112` is the shape this owes, throwing by name for anything but its one ported activation. The released `config.json` carries `"hidden_act": "silu"`, so nothing shipped is affected. Owed: parse it and refuse by name mirroring upstream's message; whether the same guard lands on the SACRED `deepseek_v2.cpp` path is a separate decision with its own red-before evidence. Found by the fresh review of [#2187](https://github.com/mudler/vllm.cpp/pull/2187) as F6. Under `## Owed` in [specs/dots3-note.md](../specs/dots3-note.md) | bug | -| [#2193](https://github.com/mudler/vllm.cpp/issues/2193) | `MODEL-MM-dots3-note` | **dots3-note's grouped MoE arm re-introduced #237's address-keyed residency and cited the #237 fix as its warrant.** W5 shipped `Dots3NoteMoePtrsFor` as `static std::map table; return table[key];` — the exact shape `ce2349dee` (2026-08-10) removed from `qwen3_5.cpp`, eighteen days earlier — under a comment claiming `qwen3_5.cpp`'s `MoeBf16Resident` relies on it. It does not: `MoeBf16ResidentFor` is `ResidentIn(w->resident_bf16)` and `ResidentIn`'s own comment says it REPLACED that map because "keying on the weight's ADDRESS let a second engine inherit a freed engine's device pointers (issue #237)". Load A, destroy it, load B at A's address: `mr.ready` is already true, the upload is skipped, and every routed expert GEMM reads A's pointers. The buffers are deliberately never freed, so there is no crash and no error — B silently answers from A's experts, quieter than the #237 repro's zeroed token ids. Secondary: `table[key]` mutates a `std::map` under no lock where `ResidentIn` takes a mutex. FIXED IN FLOW in [#2187](https://github.com/mudler/vllm.cpp/pull/2187) (fresh-review F1): `Dots3NoteMoeWeights` gained a `ResidentSlot resident_moe`, the accessor builds into it under a mutex (the `laguna.cpp:497-507` shape), and `test_moe_resident_lifetime.cpp` gained four cases for this block. What is NOT gated, said plainly: `kMoeGroupedGemmBf16` is CUDA-only with no CPU reference tier and the accessor is file-local, so no CPU gate can call it — the cases pin that residency is a member of the weights, not the accessor body. Still owed: `deepseek_v2.cpp`'s `MoePtrs` (`04f5c01e7`, 2026-07-22) carries the same pre-#237 shape and is unswept debt on a SACRED path, not touched here. Under `## Owed` in [specs/dots3-note.md](../specs/dots3-note.md) | bug | -| [#2181](https://github.com/mudler/vllm.cpp/issues/2181) | `QUANT-EXL3` | **EXL3 is not a quantization scheme in this tree, it is a DeepSeek-V4-private arm: `grep -rl Exl3 src/vllm include/vllm` returns three files, all DeepSeek-V4, and `.agents/quantization-matrix.md` carried no EXL3 row while registering 20+ other schemes.** `IsExl3Checkpoint` (`deepseek_v4_weights.cpp:229-233`) reads the same `quantization_config.quant_method == "exl3"` marker every EXL3 checkpoint carries and is consulted only from the DeepSeek-V4 loader, so no other architecture reaches the trellis kernels and no stock checkpoint loads — the parallel-path shape `AGENTS.md` §"Shared seams" forbids. **The kernels are ready and DEVICE-PROVEN**: `had_r_128` CUDA-vs-CPU `mismatches == 0`, `exl3_gemm` vs the f64 reference `rel_rms 5.538e-4` (bound `1.0e-3`), GEMV tier 3c `rel_rms 5.160e-4` (bound `6.0e-3`), on `dgx:gpu0` GB10 `sm_121a` with an `sm_121a` cubin in `cuda_exl3.cu.o`, driver 580.173.02, nvcc 13.0.88, tree `525d2b991`, 2026-08-28 — which also retires `MODEL-DSV4-EXL3` `## Owed`'s "the CUDA arm compiles nowhere yet". What is missing is everything ABOVE the kernels. **The stock layout is SIMPLER than the one implemented, measured by range-reading the safetensors header of `turboderp/Llama-3.2-1B-Instruct-exl3` @ `3.0bpw` (`f8f438c2`), 373 tensors in one 1.09 GB file**: HF-standard keys with `{trellis,suh,svh}` appended, `trellis` `I16 [k/16, n/16, 16*bits]` exactly as `Exl3ReconstructInner` reads it, norms and `embed_tokens` left `F16`, and **no `.rank{r}` segments at all** — the rank-sliced `rank-sliced-deepseek-v4-v1` schema W1b implements is SparkInfer's variant, not the format's ordinary shape. **`bits` is PER TENSOR and the config scalar is not it**: `lm_head.trellis [128, 8016, 96]` is 6-bit while `quantization_config.bits` says `3.0` and the body is 3-bit, so a reader that trusts the config decodes the head at the wrong width and no shape check catches it; the CUDA arm instantiates `bits == 3, codebook == 1` only, so that head has no device arm today. Scope in `## Scope, in waves`: W1 the `Exl3LinearMethod` on vLLM's own `LinearMethodBase` seam plus a native-layout reader and one model end to end, W2 device residency (also `MODEL-DSV4-EXL3` `## Owed`, and the precondition for the `vt::Exl3MoeMlp` device arm that skips today on `DeviceMemoryIsHostAddressable()`), W3 width coverage, W4 DeepSeek-V4 routed onto the shared seam. **The gate is the hard part and is chosen before the code**: vLLM registers no EXL3 at the parity pin and the secondary oracle does not build on aarch64 ([#1901](https://github.com/mudler/vllm.cpp/issues/1901)), so no oracle token match is available on this fleet; `## Gates` binds a bounded logit-divergence gate against the BF16 model we already gate token-exact 16/16 vs vLLM, plus a token-exact dequant-vs-native self-consistency gate, and records coherence as WEAK rather than as a pass. Spec [`quant-exl3-shared.md`](../specs/quant-exl3-shared.md) | feature | -| [#2137](https://github.com/mudler/vllm.cpp/issues/2137) | `KERNEL-ATTN-DENSE-FLASH` | **The attention-rung checker matches one spelling, so four ways of reaching `kAttention` are invisible to it — and after [#1552](https://github.com/mudler/vllm.cpp/issues/1552) that checker is the whole of the protection.** The seam decision #1552 escalated was ratified on 2026-08-27 as option (a): keep `vt::Attention` caller-opt-in, strengthened by the build-time gate, over (b) a runtime warning and (c) shape-routing `kAttention`. #1552 widened the checker's POPULATION from two non-recursive directories to `src/`, `include/` and `examples/` over every C++ suffix, closing two MEASURED holes — an unmarked call in `src/vllm/v1/attention/backend.cpp` and one in a model subdirectory each left it at `rc=0` with the OK line still reporting the same 8 sites. It did not touch the SPELLING, and the docstring has always named four that reach the same kernel undetected: `using vt::Attention;` then a bare `Attention(...)`, a `namespace vv = vt;` alias, a `#define`, and a call through a function pointer. None exists in this tree and the repository does not write attention calls that way, so this is a stated bound and not a live defect; it is filed because after (a) it is the ONLY remaining way for the #1544 failure — correct output at up to ~500x the cost, no refusal, no warning, `GetOpProviderStats` counting the naive selection as the success it genuinely is — to recur with nothing firing. **Widening the regex is not the repair, and that is pinned rather than asserted:** `\bAttention\s*\(` also matches every fast rung's suffix-free form and would demand a marker beside exactly the calls the checker wants people to make, which `test_check_attention_rung_consistency.py::test_widening_the_regex_to_the_fast_rungs_is_visible` already holds, and no regex reaches a function pointer at all. What closes it is a compiler-side population — the CUDA op registry, or a clang tooling pass over the real translation unit — which is a different instrument and its own row. So a green means "no unmarked `vt::Attention(` call in the scanned population", never "no model is on the naive rung", which is the honest limit of the ratified seam. NOT fixed in flow: it needs an instrument this tree does not have. Owner: row `KERNEL-ATTN-DENSE-FLASH`, under `## Owed` in [`eng-attn-optin-sweep.md`](../specs/eng-attn-optin-sweep.md) | gap | -| [#2140](https://github.com/mudler/vllm.cpp/issues/2140) | `LTX25-TEXT-PROJ-DTYPE` | **The LTX-2.5 caption projections are NVFP4-only, so the bf16 text tower loads and the render still refuses.** `LoadProjection` (`src/vllm/model_executor/models/ltx2_loader.cpp`) computes `in_features = w->shape[1] * 2` unconditionally, with the comment "NVFP4 packs TWO values per byte", then requires `.weight_scale` and `.weight_scale_2` and dequantizes. On the bf16 checkpoint the stored width is already logical, so the doubling turns a correct 188160 into 376320 and the geometry check fires on the loader's own arithmetic. Measured on `dgx:gpu0` (`rc` job `001c36e9-76b1-432c-9536-2d24c0e613d0`, 2026-08-27) and confirmed by reading both safetensors headers: the bf16 file stores `text_embedding_projection.video_aggregate_embed.weight` as `BF16 [4096, 188160]` with **zero** `.weight_scale` tensors and **zero** `torchao_nvfp4` markers in the whole file, while the torchao file stores it as `U8 [4096, 94080]` with 334 of each. The fix resolves the storage format from the file the way upstream does — `_discover_nvfp4_layers` (`packages/ltx-core/src/ltx_core/quantization/nvfp4/prequant.py:30-50` at pin `fd4ded7f`) selects a layer only when `.weight_scale` and `.weight_scale_2` are BOTH present and the dtype triple is `U8`/`F8_E4M3`/`F32`, treats exactly one of the pair as an error, and leaves everything else the plain `nn.Linear(flat_dim, ...)` of `encoder_configurator.py:206-208`, whose stored width IS its logical width. Blocks [#1854](https://github.com/mudler/vllm.cpp/issues/1854)'s absolute gate, because [#1864](https://github.com/mudler/vllm.cpp/issues/1864)'s reference render was taken with the bf16 tower and an arm-matched comparison cannot substitute the NVFP4 one. Spec [`ltx25-text-proj-dtype.md`](../specs/ltx25-text-proj-dtype.md) | bug | -| [#2130](https://github.com/mudler/vllm.cpp/issues/2130) | `LTX25-ORACLE-ABSOLUTE` | **`vllm_video_params.steps` ships in the ABI, the engine honours it, and no shipped client can set it, so every LTX-2.5 render silently runs the recipe default.** `include/vllm.h:1075` declares `int32_t steps`, `src/capi/vllm_c.cpp:1664` forwards it, and `src/vllm/multimodal/ltx2_video.cpp:4027` reads it — `int64_t steps = gen.steps > 0 ? gen.steps : recipe.num_inference_steps;` — with `allow_request_sigmas` true and `fixed_num_inference_steps` false (`include/vllm/model_executor/models/ltx2_pipeline.h:872-877`), so a request value is honoured rather than clamped. `examples/ltx2_gen/main.cpp:306-451` parses `--frames`, `--width`, `--height` and `--seed` and assigns `steps` nowhere, so `one_stage` at model version 2.5 always runs **30** (`ltx2_pipeline.cpp:1157` from `Ltx2Params24()` → `Ltx2Params23()`, where `:968` sets `num_inference_steps = 30`). WHAT IT COST: [#1864](https://github.com/mudler/vllm.cpp/issues/1864)'s reference render was taken at **8** steps (`tools/oracle/ltx2_oracle.py:88`), so [#1854](https://github.com/mudler/vllm.cpp/issues/1854)'s absolute comparison had a 3.75x denoise-budget confound on the one axis the CLI cannot reach — and it confounds in the direction that FLATTERS us, so a pass taken on it would be unearned. The reference is arm-matched on all four checkpoints, geometry, seed and prompt; this was the only unmatched axis. The `AGENTS.md` "Nothing lands dead" shape at the SEAM rather than in the engine: the capability is reachable through `include/vllm.h` and the thin ABI client that exposes every neighbouring field does not expose this one, and nothing detects it because the renders are correct, no refusal fires, and the only symptom is that every LTX-2.5 render in this tree has run one step count. FIXED IN FLOW: `--steps N` forwards to the existing ABI field and to nothing else; no engine code changes. Owner: row `LTX25-ORACLE-ABSOLUTE`, spec [`ltx25-oracle-absolute.md`](../specs/ltx25-oracle-absolute.md) | bug | -| [#2157](https://github.com/mudler/vllm.cpp/issues/2157) | — | `scripts/check-commit-trailers.py:463` walks `rev-list --reverse base..head` with no `--no-merges` and no parent-count test, so a plain `git merge origin/main` on a task branch reds `commit-protocol-tag` with three findings per merge commit — while the SAME job skips merge commits 50 lines earlier (`.github/workflows/ci.yml:873`, "they are not authored content"), so one job carries two opposite rules and only one of them is written down. Measured: `--range a0f12b727..d05723f8e` is rc=1 with 9 findings across 3 merge OIDs, while the same range's 3 NON-merge commits all pass, which is the isolating control; CI agrees on PR #2134 (job 98706339787) over `0d8962500cc1`, two parents and a 0-byte body. Nothing reaches `main`: `squash_merge_commit_message = PR_BODY` means a branch merge commit never becomes a landed message, so the cost is a red gate plus a forced branch rewrite on every branch that syncs — which AGENTS.md § Landing work instructs as the routine response to a rejected push. AGENTS.md is SILENT on merge commits (`grep -rn 'merges included'` returns nothing against a positive control returning 33), so this is a gap rather than a policy. NOT FIXED: changing the walk is a semantic gate change owing its own row, spec, red-before and green-after, and the developer chose on 2026-08-28 to authorize `row/*` force-push instead. #1136 (CLOSED) records the same mechanism as one PR's review finding and owns no repair; #581 is the forge's merges on `main`; #467 is preflight not running the checker; #406 is trailer-block LOCATION and leaves this shape red on purpose. Owed under `## Owed` in [`fix-trailer-lane-cutover.md`](../specs/fix-trailer-lane-cutover.md) | bug | -| [#2184](https://github.com/mudler/vllm.cpp/issues/2184) | `FIX-TRAILER-LANE-CUTOVER` | Four `ci.yml` line anchors in [`fix-trailer-lane-cutover.md`](../specs/fix-trailer-lane-cutover.md) resolved to code that does not support the sentence citing them, because the strict trailer walk MOVED from `agent-record` (#863) to `commit-protocol-tag` (`ci.yml:818`) after the spec was written: the prose survived the move and the numbers did not. `:596-623` landed on a GPU-mutex comment and `pending_args`; `:626-635` on a bare `fi` and a `--pr-base` continuation. NOT a guess to repair — spec line 44 names its own job in the sentence ("in `commit-protocol-tag`"), so it is right-job/wrong-number, and the other two claims each map to a unique construct. Repointed at `6f02680bb` to `:899-927` (LAST_GREEN `:899`, base `:917-920`, walk `:927`), `:935-939` (the `--filled` body guard), `:924`/`:927` (the only two `--range`-alone calls) and `agent-integration.py:106-110` (tightened; `--cutover` is on `:108`). No checker can see this class: `check-symbol-anchors` resolves SYMBOLS, not whether a line range supports a claim, so the pointer lands on plausible code and the reader finds nothing to contradict them. Fixed in flow in PR #2159 | bug | +| [#2190](https://github.com/mudler/vllm.cpp/issues/2190) | `MODEL-MM-dots3-note` | **dots3-note's quantization refusal keys on `weight_block_size` alone, so a per-tensor or per-row fp8/gptq/awq config loads and silently dequantizes to bf16.** `Dots3NoteDeviceRefusal` (`dots3_note_device.cpp:855`) branches on `has_blockwise_quant()`, which is `!weight_block_size.empty()` (`dots3_note.h:208`). `quant_method` IS parsed (`dots3_note.cpp:264-268`) and stored (`dots3_note.h:206`), and is read for nothing but the text of the blockwise message (`:862-863`). A `config.json` with `quantization_config.quant_method = "fp8"` (or gptq/awq/mxfp4/compressed-tensors) and NO `weight_block_size` therefore passes, and `dense_loaders::MaterializeBf16Source` silently dequantizes a per-tensor or per-output-ROW `_scale` into a bf16 GEMM — which is precisely the case the refusal's own comment names as the worse one, five lines above the branch that does not cover it (`:849-854`). This row has NO oracle on any hardware we own (spec §6.4), so nothing downstream catches the plausible wrong answer. No released checkpoint is affected: the bf16 repo carries no `quantization_config` and the `-fp8` sibling carries `weight_block_size [128, 128]` and is refused correctly. Owed: refuse a non-empty `quant_method` this port cannot read, naming the method and W9, with the config-fixture gate the blockwise case already has. Found by the fresh review of [#2187](https://github.com/mudler/vllm.cpp/pull/2187) as F5 and deliberately not fixed there — a refusal-semantics change needs its own red-before fixture. Under `## Owed` in [specs/dots3-note.md](specs/dots3-note.md) | bug | +| [#2191](https://github.com/mudler/vllm.cpp/issues/2191) | `MODEL-MM-dots3-note` | **`hidden_act` is not mirrored, so a non-silu dots3-note config runs SwiGLU silently where vLLM raises.** `DeepseekV2MoE.__init__` refuses anything but silu before it builds a thing — `deepseek_v2.py:310-314` @ pin `bc2d63e650`, `ValueError(f"Unsupported activation: {config.hidden_act}. Only silu is supported for now.")` — and `Dots3NoteLanguageModelForCausalLM` (`model.py:681`) subclasses `DeepseekV32ForCausalLM`, so that is the `__init__` W5 ports. `grep -c hidden_act` over `dots3_note.cpp` and `dots3_note.h` is **0**: the key is never parsed, `Dots3NoteParams` has no field for it, and `Dots3NoteDeviceRefusal` never mentions it, so `hidden_act: "gelu"` loads and runs `vt::MoeSiluMul` / `vt::MoeGroupedGemmBf16GateUpSilu` with no refusal. Same hole in `deepseek_v2.cpp` and `deepseek_v2_weights.cpp` (`grep -c` = 0 on both), so it is a MIRROR GAP inherited by both ports rather than a W5 regression; `parakeet_transducer.cpp:108-112` is the shape this owes, throwing by name for anything but its one ported activation. The released `config.json` carries `"hidden_act": "silu"`, so nothing shipped is affected. Owed: parse it and refuse by name mirroring upstream's message; whether the same guard lands on the SACRED `deepseek_v2.cpp` path is a separate decision with its own red-before evidence. Found by the fresh review of [#2187](https://github.com/mudler/vllm.cpp/pull/2187) as F6. Under `## Owed` in [specs/dots3-note.md](specs/dots3-note.md) | bug | +| [#2193](https://github.com/mudler/vllm.cpp/issues/2193) | `MODEL-MM-dots3-note` | **dots3-note's grouped MoE arm re-introduced #237's address-keyed residency and cited the #237 fix as its warrant.** W5 shipped `Dots3NoteMoePtrsFor` as `static std::map table; return table[key];` — the exact shape `ce2349dee` (2026-08-10) removed from `qwen3_5.cpp`, eighteen days earlier — under a comment claiming `qwen3_5.cpp`'s `MoeBf16Resident` relies on it. It does not: `MoeBf16ResidentFor` is `ResidentIn(w->resident_bf16)` and `ResidentIn`'s own comment says it REPLACED that map because "keying on the weight's ADDRESS let a second engine inherit a freed engine's device pointers (issue #237)". Load A, destroy it, load B at A's address: `mr.ready` is already true, the upload is skipped, and every routed expert GEMM reads A's pointers. The buffers are deliberately never freed, so there is no crash and no error — B silently answers from A's experts, quieter than the #237 repro's zeroed token ids. Secondary: `table[key]` mutates a `std::map` under no lock where `ResidentIn` takes a mutex. FIXED IN FLOW in [#2187](https://github.com/mudler/vllm.cpp/pull/2187) (fresh-review F1): `Dots3NoteMoeWeights` gained a `ResidentSlot resident_moe`, the accessor builds into it under a mutex (the `laguna.cpp:497-507` shape), and `test_moe_resident_lifetime.cpp` gained four cases for this block. What is NOT gated, said plainly: `kMoeGroupedGemmBf16` is CUDA-only with no CPU reference tier and the accessor is file-local, so no CPU gate can call it — the cases pin that residency is a member of the weights, not the accessor body. Still owed: `deepseek_v2.cpp`'s `MoePtrs` (`04f5c01e7`, 2026-07-22) carries the same pre-#237 shape and is unswept debt on a SACRED path, not touched here. Under `## Owed` in [specs/dots3-note.md](specs/dots3-note.md) | bug | +| [#2181](https://github.com/mudler/vllm.cpp/issues/2181) | `QUANT-EXL3` | **EXL3 is not a quantization scheme in this tree, it is a DeepSeek-V4-private arm: `grep -rl Exl3 src/vllm include/vllm` returns three files, all DeepSeek-V4, and `.agents/quantization-matrix.md` carried no EXL3 row while registering 20+ other schemes.** `IsExl3Checkpoint` (`deepseek_v4_weights.cpp:229-233`) reads the same `quantization_config.quant_method == "exl3"` marker every EXL3 checkpoint carries and is consulted only from the DeepSeek-V4 loader, so no other architecture reaches the trellis kernels and no stock checkpoint loads — the parallel-path shape `AGENTS.md` §"Shared seams" forbids. **The kernels are ready and DEVICE-PROVEN**: `had_r_128` CUDA-vs-CPU `mismatches == 0`, `exl3_gemm` vs the f64 reference `rel_rms 5.538e-4` (bound `1.0e-3`), GEMV tier 3c `rel_rms 5.160e-4` (bound `6.0e-3`), on `dgx:gpu0` GB10 `sm_121a` with an `sm_121a` cubin in `cuda_exl3.cu.o`, driver 580.173.02, nvcc 13.0.88, tree `525d2b991`, 2026-08-28 — which also retires `MODEL-DSV4-EXL3` `## Owed`'s "the CUDA arm compiles nowhere yet". What is missing is everything ABOVE the kernels. **The stock layout is SIMPLER than the one implemented, measured by range-reading the safetensors header of `turboderp/Llama-3.2-1B-Instruct-exl3` @ `3.0bpw` (`f8f438c2`), 373 tensors in one 1.09 GB file**: HF-standard keys with `{trellis,suh,svh}` appended, `trellis` `I16 [k/16, n/16, 16*bits]` exactly as `Exl3ReconstructInner` reads it, norms and `embed_tokens` left `F16`, and **no `.rank{r}` segments at all** — the rank-sliced `rank-sliced-deepseek-v4-v1` schema W1b implements is SparkInfer's variant, not the format's ordinary shape. **`bits` is PER TENSOR and the config scalar is not it**: `lm_head.trellis [128, 8016, 96]` is 6-bit while `quantization_config.bits` says `3.0` and the body is 3-bit, so a reader that trusts the config decodes the head at the wrong width and no shape check catches it; the CUDA arm instantiates `bits == 3, codebook == 1` only, so that head has no device arm today. Scope in `## Scope, in waves`: W1 the `Exl3LinearMethod` on vLLM's own `LinearMethodBase` seam plus a native-layout reader and one model end to end, W2 device residency (also `MODEL-DSV4-EXL3` `## Owed`, and the precondition for the `vt::Exl3MoeMlp` device arm that skips today on `DeviceMemoryIsHostAddressable()`), W3 width coverage, W4 DeepSeek-V4 routed onto the shared seam. **The gate is the hard part and is chosen before the code**: vLLM registers no EXL3 at the parity pin and the secondary oracle does not build on aarch64 ([#1901](https://github.com/mudler/vllm.cpp/issues/1901)), so no oracle token match is available on this fleet; `## Gates` binds a bounded logit-divergence gate against the BF16 model we already gate token-exact 16/16 vs vLLM, plus a token-exact dequant-vs-native self-consistency gate, and records coherence as WEAK rather than as a pass. Spec [`quant-exl3-shared.md`](specs/quant-exl3-shared.md) | feature | +| [#2137](https://github.com/mudler/vllm.cpp/issues/2137) | `KERNEL-ATTN-DENSE-FLASH` | **The attention-rung checker matches one spelling, so four ways of reaching `kAttention` are invisible to it — and after [#1552](https://github.com/mudler/vllm.cpp/issues/1552) that checker is the whole of the protection.** The seam decision #1552 escalated was ratified on 2026-08-27 as option (a): keep `vt::Attention` caller-opt-in, strengthened by the build-time gate, over (b) a runtime warning and (c) shape-routing `kAttention`. #1552 widened the checker's POPULATION from two non-recursive directories to `src/`, `include/` and `examples/` over every C++ suffix, closing two MEASURED holes — an unmarked call in `src/vllm/v1/attention/backend.cpp` and one in a model subdirectory each left it at `rc=0` with the OK line still reporting the same 8 sites. It did not touch the SPELLING, and the docstring has always named four that reach the same kernel undetected: `using vt::Attention;` then a bare `Attention(...)`, a `namespace vv = vt;` alias, a `#define`, and a call through a function pointer. None exists in this tree and the repository does not write attention calls that way, so this is a stated bound and not a live defect; it is filed because after (a) it is the ONLY remaining way for the #1544 failure — correct output at up to ~500x the cost, no refusal, no warning, `GetOpProviderStats` counting the naive selection as the success it genuinely is — to recur with nothing firing. **Widening the regex is not the repair, and that is pinned rather than asserted:** `\bAttention\s*\(` also matches every fast rung's suffix-free form and would demand a marker beside exactly the calls the checker wants people to make, which `test_check_attention_rung_consistency.py::test_widening_the_regex_to_the_fast_rungs_is_visible` already holds, and no regex reaches a function pointer at all. What closes it is a compiler-side population — the CUDA op registry, or a clang tooling pass over the real translation unit — which is a different instrument and its own row. So a green means "no unmarked `vt::Attention(` call in the scanned population", never "no model is on the naive rung", which is the honest limit of the ratified seam. NOT fixed in flow: it needs an instrument this tree does not have. Owner: row `KERNEL-ATTN-DENSE-FLASH`, under `## Owed` in [`eng-attn-optin-sweep.md`](specs/eng-attn-optin-sweep.md) | gap | +| [#2140](https://github.com/mudler/vllm.cpp/issues/2140) | `LTX25-TEXT-PROJ-DTYPE` | **The LTX-2.5 caption projections are NVFP4-only, so the bf16 text tower loads and the render still refuses.** `LoadProjection` (`src/vllm/model_executor/models/ltx2_loader.cpp`) computes `in_features = w->shape[1] * 2` unconditionally, with the comment "NVFP4 packs TWO values per byte", then requires `.weight_scale` and `.weight_scale_2` and dequantizes. On the bf16 checkpoint the stored width is already logical, so the doubling turns a correct 188160 into 376320 and the geometry check fires on the loader's own arithmetic. Measured on `dgx:gpu0` (`rc` job `001c36e9-76b1-432c-9536-2d24c0e613d0`, 2026-08-27) and confirmed by reading both safetensors headers: the bf16 file stores `text_embedding_projection.video_aggregate_embed.weight` as `BF16 [4096, 188160]` with **zero** `.weight_scale` tensors and **zero** `torchao_nvfp4` markers in the whole file, while the torchao file stores it as `U8 [4096, 94080]` with 334 of each. The fix resolves the storage format from the file the way upstream does — `_discover_nvfp4_layers` (`packages/ltx-core/src/ltx_core/quantization/nvfp4/prequant.py:30-50` at pin `fd4ded7f`) selects a layer only when `.weight_scale` and `.weight_scale_2` are BOTH present and the dtype triple is `U8`/`F8_E4M3`/`F32`, treats exactly one of the pair as an error, and leaves everything else the plain `nn.Linear(flat_dim, ...)` of `encoder_configurator.py:206-208`, whose stored width IS its logical width. Blocks [#1854](https://github.com/mudler/vllm.cpp/issues/1854)'s absolute gate, because [#1864](https://github.com/mudler/vllm.cpp/issues/1864)'s reference render was taken with the bf16 tower and an arm-matched comparison cannot substitute the NVFP4 one. Spec [`ltx25-text-proj-dtype.md`](specs/ltx25-text-proj-dtype.md) | bug | +| [#2130](https://github.com/mudler/vllm.cpp/issues/2130) | `LTX25-ORACLE-ABSOLUTE` | **`vllm_video_params.steps` ships in the ABI, the engine honours it, and no shipped client can set it, so every LTX-2.5 render silently runs the recipe default.** `include/vllm.h:1075` declares `int32_t steps`, `src/capi/vllm_c.cpp:1664` forwards it, and `src/vllm/multimodal/ltx2_video.cpp:4027` reads it — `int64_t steps = gen.steps > 0 ? gen.steps : recipe.num_inference_steps;` — with `allow_request_sigmas` true and `fixed_num_inference_steps` false (`include/vllm/model_executor/models/ltx2_pipeline.h:872-877`), so a request value is honoured rather than clamped. `examples/ltx2_gen/main.cpp:306-451` parses `--frames`, `--width`, `--height` and `--seed` and assigns `steps` nowhere, so `one_stage` at model version 2.5 always runs **30** (`ltx2_pipeline.cpp:1157` from `Ltx2Params24()` → `Ltx2Params23()`, where `:968` sets `num_inference_steps = 30`). WHAT IT COST: [#1864](https://github.com/mudler/vllm.cpp/issues/1864)'s reference render was taken at **8** steps (`tools/oracle/ltx2_oracle.py:88`), so [#1854](https://github.com/mudler/vllm.cpp/issues/1854)'s absolute comparison had a 3.75x denoise-budget confound on the one axis the CLI cannot reach — and it confounds in the direction that FLATTERS us, so a pass taken on it would be unearned. The reference is arm-matched on all four checkpoints, geometry, seed and prompt; this was the only unmatched axis. The `AGENTS.md` "Nothing lands dead" shape at the SEAM rather than in the engine: the capability is reachable through `include/vllm.h` and the thin ABI client that exposes every neighbouring field does not expose this one, and nothing detects it because the renders are correct, no refusal fires, and the only symptom is that every LTX-2.5 render in this tree has run one step count. FIXED IN FLOW: `--steps N` forwards to the existing ABI field and to nothing else; no engine code changes. Owner: row `LTX25-ORACLE-ABSOLUTE`, spec [`ltx25-oracle-absolute.md`](specs/ltx25-oracle-absolute.md) | bug | +| [#2157](https://github.com/mudler/vllm.cpp/issues/2157) | — | `scripts/check-commit-trailers.py:463` walks `rev-list --reverse base..head` with no `--no-merges` and no parent-count test, so a plain `git merge origin/main` on a task branch reds `commit-protocol-tag` with three findings per merge commit — while the SAME job skips merge commits 50 lines earlier (`.github/workflows/ci.yml:873`, "they are not authored content"), so one job carries two opposite rules and only one of them is written down. Measured: `--range a0f12b727..d05723f8e` is rc=1 with 9 findings across 3 merge OIDs, while the same range's 3 NON-merge commits all pass, which is the isolating control; CI agrees on PR #2134 (job 98706339787) over `0d8962500cc1`, two parents and a 0-byte body. Nothing reaches `main`: `squash_merge_commit_message = PR_BODY` means a branch merge commit never becomes a landed message, so the cost is a red gate plus a forced branch rewrite on every branch that syncs — which AGENTS.md § Landing work instructs as the routine response to a rejected push. AGENTS.md is SILENT on merge commits (`grep -rn 'merges included'` returns nothing against a positive control returning 33), so this is a gap rather than a policy. NOT FIXED: changing the walk is a semantic gate change owing its own row, spec, red-before and green-after, and the developer chose on 2026-08-28 to authorize `row/*` force-push instead. #1136 (CLOSED) records the same mechanism as one PR's review finding and owns no repair; #581 is the forge's merges on `main`; #467 is preflight not running the checker; #406 is trailer-block LOCATION and leaves this shape red on purpose. Owed under `## Owed` in [`fix-trailer-lane-cutover.md`](specs/fix-trailer-lane-cutover.md) | bug | +| [#2184](https://github.com/mudler/vllm.cpp/issues/2184) | `FIX-TRAILER-LANE-CUTOVER` | Four `ci.yml` line anchors in [`fix-trailer-lane-cutover.md`](specs/fix-trailer-lane-cutover.md) resolved to code that does not support the sentence citing them, because the strict trailer walk MOVED from `agent-record` (#863) to `commit-protocol-tag` (`ci.yml:818`) after the spec was written: the prose survived the move and the numbers did not. `:596-623` landed on a GPU-mutex comment and `pending_args`; `:626-635` on a bare `fi` and a `--pr-base` continuation. NOT a guess to repair — spec line 44 names its own job in the sentence ("in `commit-protocol-tag`"), so it is right-job/wrong-number, and the other two claims each map to a unique construct. Repointed at `6f02680bb` to `:899-927` (LAST_GREEN `:899`, base `:917-920`, walk `:927`), `:935-939` (the `--filled` body guard), `:924`/`:927` (the only two `--range`-alone calls) and `agent-integration.py:106-110` (tightened; `--cutover` is on `:108`). No checker can see this class: `check-symbol-anchors` resolves SYMBOLS, not whether a line range supports a claim, so the pointer lands on plausible code and the reader finds nothing to contradict them. Fixed in flow in PR #2159 | bug | | [#2098](https://github.com/mudler/vllm.cpp/issues/2098) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash's mHC head collapse is an unweighted mean, and DeepSeek-V4's gated `HcHeadCollapse` is the wrong final projection.** W4 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998). `Glm5NextTextHyperHead.forward` is `hidden_streams.mean(dim=2)` and its own docstring says "Unlike DeepSeek-V4" (`modular_glm5_next.py:368-372` @ transformers v5.16.1); the checkpoint carries no `hc_head.*` tensor at any layer, so there are no weights a gated collapse could read. The other three mHC pieces ARE V4's and are reused. Landed `src/vllm/model_executor/models/glm5_next_mhc.{h,cpp}` gated against goldens RUN out of the pinned reference | feature | -| [#2194](https://github.com/mudler/vllm.cpp/issues/2194) | `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` | **Both upstream anchors on the `GlmMoeDsaForCausalLM` row were stale at our own parity pin, and one of them confirmed itself to a casual reader.** At `5559679229bc961848b121ccdeaa8fa5d79bec98` the row's `registry.py:116` is `"Glm4MoeLiteForCausalLM"`, a DIFFERENT model, where the entry for this architecture is `:117`; and its `deepseek_v2.py:1917-1918` is `load_weights` / `loader = AutoWeightsLoader(self)`, where the class is `:1930`. Neither number was wrong when written: both are exact at the PRIOR pin `e24d1b24`, which the spike's `### Upstream chain` names and still names honestly. The 2026-07-26 advance moved `registry.py` by one line and `deepseek_v2.py` by thirteen, and the row carried the coordinates forward with no revision label, so nothing could see them drift. Reconciled in flow against the published `zai-org/GLM-5.3` (revision `935644c05e76`, `model_type: glm_moe_dsa`, 753,329,940,480 parameters, 703.74 GiB of fp8 over 141 shards, 1403.2 GiB at bf16), which the row predated and named nowhere. Three further corrections: "GLM-5.x is DeepSeek-V3.2 VERBATIM" now says AT THE PIN and names the three places vLLM `main` `d1922cb5a7` diverges (the alias re-homed to `vllm.models.deepseek_v32`, its own `VerifyAndUpdateConfig` at `config.py:43` registered `:936`, and membership of `DEFAULT_BREAKABLE_CUDAGRAPH_ARCHITECTURES` at `vllm/config/vllm.py:81`); the blocker arithmetic is recorded so nobody redoes it (1.3641 bpw to fit 119.631 GiB on `dgx:gpu0`, against 203.5 / 149.1 / 131.5 GiB at 2.32 / 1.70 / 1.50 bpw, and `unsloth/GLM-5.3-GGUF`'s one complete arm `UD-Q3_K_XL` at 319.41 GiB); and both oracles are registered as reaching the architecture at revisions this tree ALREADY pins — vLLM primary at `registry.py:117` and `deepseek_v2.py:1930`, llama.cpp at stock `b10451` where `LLM_ARCH_GLM_DSA` -> `"glm-dsa"` sits at `src/llama-arch.cpp:85` with its case at `:1051`, its graph at `src/models/glm-dsa.cpp` and its converter at `conversion/glm.py:274-276`. So NEITHER needs a new oracle file, unlike `llama-cpp-qwen4exp` and the `llama-cpp-glm5next` proposed in [#2178](https://github.com/mudler/vllm.cpp/issues/2178), which exist only because no llama.cpp RELEASE carries those architectures. Both are `gateable = no` FOR THIS MODEL on MEMORY, not on missing support, which is what separates this row from `MODEL-MM-GLM53-FLASH`. The row stays `BLOCKED`; records only, no product code, no pin advance, no second matrix row. Spec [`glm-dsa-latest-deepseek.md`](../specs/glm-dsa-latest-deepseek.md) §2 | bug | +| [#2194](https://github.com/mudler/vllm.cpp/issues/2194) | `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` | **Both upstream anchors on the `GlmMoeDsaForCausalLM` row were stale at our own parity pin, and one of them confirmed itself to a casual reader.** At `5559679229bc961848b121ccdeaa8fa5d79bec98` the row's `registry.py:116` is `"Glm4MoeLiteForCausalLM"`, a DIFFERENT model, where the entry for this architecture is `:117`; and its `deepseek_v2.py:1917-1918` is `load_weights` / `loader = AutoWeightsLoader(self)`, where the class is `:1930`. Neither number was wrong when written: both are exact at the PRIOR pin `e24d1b24`, which the spike's `### Upstream chain` names and still names honestly. The 2026-07-26 advance moved `registry.py` by one line and `deepseek_v2.py` by thirteen, and the row carried the coordinates forward with no revision label, so nothing could see them drift. Reconciled in flow against the published `zai-org/GLM-5.3` (revision `935644c05e76`, `model_type: glm_moe_dsa`, 753,329,940,480 parameters, 703.74 GiB of fp8 over 141 shards, 1403.2 GiB at bf16), which the row predated and named nowhere. Three further corrections: "GLM-5.x is DeepSeek-V3.2 VERBATIM" now says AT THE PIN and names the three places vLLM `main` `d1922cb5a7` diverges (the alias re-homed to `vllm.models.deepseek_v32`, its own `VerifyAndUpdateConfig` at `config.py:43` registered `:936`, and membership of `DEFAULT_BREAKABLE_CUDAGRAPH_ARCHITECTURES` at `vllm/config/vllm.py:81`); the blocker arithmetic is recorded so nobody redoes it (1.3641 bpw to fit 119.631 GiB on `dgx:gpu0`, against 203.5 / 149.1 / 131.5 GiB at 2.32 / 1.70 / 1.50 bpw, and `unsloth/GLM-5.3-GGUF`'s one complete arm `UD-Q3_K_XL` at 319.41 GiB); and both oracles are registered as reaching the architecture at revisions this tree ALREADY pins — vLLM primary at `registry.py:117` and `deepseek_v2.py:1930`, llama.cpp at stock `b10451` where `LLM_ARCH_GLM_DSA` -> `"glm-dsa"` sits at `src/llama-arch.cpp:85` with its case at `:1051`, its graph at `src/models/glm-dsa.cpp` and its converter at `conversion/glm.py:274-276`. So NEITHER needs a new oracle file, unlike `llama-cpp-qwen4exp` and the `llama-cpp-glm5next` proposed in [#2178](https://github.com/mudler/vllm.cpp/issues/2178), which exist only because no llama.cpp RELEASE carries those architectures. Both are `gateable = no` FOR THIS MODEL on MEMORY, not on missing support, which is what separates this row from `MODEL-MM-GLM53-FLASH`. The row stays `BLOCKED`; records only, no product code, no pin advance, no second matrix row. Spec [`glm-dsa-latest-deepseek.md`](specs/glm-dsa-latest-deepseek.md) §2 | bug | | [#2199](https://github.com/mudler/vllm.cpp/issues/2199) | `BACKEND-TENSTORRENT-QWEN35` | **The row spec's `## Now` still says "Owed next: W4 — cut the host staging wall (this row's active gate)" after W4 landed.** #2118 landed levers 1+2 on 2026-08-28 (`7ba0dfe1a`: bulk bf16 staging + single-slot resolution, 0.104 → 0.177 tok/s, `Numel()` 27.09% → 1.76%, review PASS in `f99116ce2`), and #2115's opt-out-arm pair (`3fe34e2c6`) landed after that, but the section was written before W4 and neither landing carried the reconciliation the section itself scheduled ("Before W4, reconcile … `## Git integration`'s base, and this section itself"). `scripts/now.py` renders the row's live Next step from this section, so the derived surface reports an already-landed wave as the active gate. Record-only repair: rewrite `## Now` to the post-#2118 position, bump `## Git integration`'s base `8f5d4e4ed` → `3fe34e2c6`, drop the now-landed USAGE.md weights clause | bug | | [#2211](https://github.com/mudler/vllm.cpp/issues/2211) | `MODEL-MM-QWEN4-EXP` | **The QSA indexer composition lived in a TEST HELPER, so nothing under `src/` enforced any of the four settings it depends on.** W5b-5 of [#2031](https://github.com/mudler/vllm.cpp/issues/2031). W5b-4 landed the claim that QSA's block score and top-k ARE `vt::DsaIndexerLogits` + `vt::DsaTopkSelect` with the per-head fold collapsed, and that collapse holds only under four settings — `weights` all ones, `n_head_scale == 1` (not DeepSeek-V4's `n_head ** -0.5`), `softmax_scale == index_head_dim ** -0.5`, and `win_end == kv_len / compress_ratio` per query token. All four lived in `RunIndexer` in `tests/vllm/models/test_qwen4_exp_qsa_device.cpp`, and TWO of them are invisible to any selection-based gate BY CONSTRUCTION: top-k is invariant under a positive rescale of every score, which is why spec mutation M26 SURVIVED. **Fixed by `src/vllm/model_executor/models/qwen4_exp_qsa_block.{h,cpp}`:** `Qwen4ExpQsaIndex` composes the three ops with the four settings stated and asserted at one site, and `RunQwen4ExpQsaBlock` is `Qwen4ExpTextAttention.forward` end to end — the q/k/v/o projections, the `(1 + w)`-polarity q/k norms (the GGUF loader INVERTS the converter's baked `+1`, so the fold belongs at the norm, and a port that passed the gamma through an `out * w` norm would apply a near-zero scale), the RoPE, the gather consumer, the sigmoid OUTPUT GATE and `o_proj`. **The VALUE gate the spec demanded now exists:** `tests/vllm/models/fixtures/gen_qwen4_exp_qsa_block_goldens.py` captures the oracle's OWN pre-top-k `scores` tensor by intercepting `torch.Tensor.topk` inside the unmodified `Qwen4ExpTextQSAIndexer.forward` at the lane pin (transformers 5.16.0), and fed the oracle's own roped query and raw keys the composed logits are BIT-IDENTICAL to it (max abs 0 over a scale of 3.37 and 6.24, 12 and 60 logits). M1 (inherit `n_head ** -0.5`) and M2 (drop the softmax scale) both RED on it, which is the repair for M26. Also gated: a released-config case at kv_len 3002, past the 2048 budget below which every candidate is selected (`keys_visited` 16400 against a dense 24016), and a NaN-poison probe at the BLOCK layer that convicts a mask-shaped consumer. 24 mutations, no survivors after one repair — writing the indexer key at row 0 instead of `past_len` SURVIVED the first battery and is closed by a row-for-row side-cache comparison after a split prefill/decode. **Lands UNREACHED**, and `## Owed` records it: `ModelRegistry::Forward` still refuses `qwen4_exp` by name because the layer loop and the KV-cache spec are not written. | feature | | [#2198](https://github.com/mudler/vllm.cpp/issues/2198) | `MODEL-MM-QWEN4-EXP` | **W4's QSA comments cited `tokens_per_state`, a field with ZERO hits over the pinned vLLM tree, and the wave writing the KV-cache spec is exactly who would have gone looking for it.** Fixed IN FLOW under W5c-1 of [#2031](https://github.com/mudler/vllm.cpp/issues/2031). `grep -rn tokens_per_state` over `/home/mudler/_git/vllm/vllm/` at the parity pin `5559679229` returns nothing tree-wide, and neither does a search for the docstring the comments quoted ("Ints > 1 compress multiple tokens into one state"); the anchor they cited, `v1/attention/backends/mla/indexer.py:624-628`, is `_prepare_decode_tensors` and is unrelated to KV sizing. The real field is **`compress_ratio`** — `vllm/v1/kv_cache_interface.py:386` declares it defaulted to 1, `:393-395` is `storage_block_size = block_size // compress_ratio`, `:617` and `:624-625` repeat the pair on `SlidingWindowMLASpec`, and `:424-435` is `MLAAttentionSpec.merge` asserting ONE `compress_ratio` per KV group. This tree was already correct where it matters (`include/vllm/v1/kv_cache_interface.h` spells it `compress_ratio`), so the defect was a CITATION and never a number: the two sites are `src/vllm/model_executor/models/qwen4_exp_qsa.h`'s port-map comment and its `QsaSideCacheSpec` doc comment, both of which now cite `compress_ratio` with the three anchors above and record what was wrong so the correction is not re-derived. `QsaSideCacheSpec::tokens_per_state` KEEPS its name deliberately — it is a LOCAL field with no upstream referent whose arithmetic is right (64 B/token/layer at bf16, pinned by `tests/vllm/models/test_qwen4_exp_qsa.cpp`) and identical to `MLAAttentionSpec::real_page_size_bytes()`, so renaming it would churn W4's TU and suite to fix a citation the comments now carry; a comment beside the field says it has no upstream referent. Found while scoping W5c, whose `MLAAttentionSpec` third group is built with `compress_ratio=4` and whose `block_size % compress_ratio` refusal exists because `storage_block_size()` truncates in silence | bug | -| [#2203](https://github.com/mudler/vllm.cpp/issues/2203) | `MODEL-MM-QWEN4-EXP` | **`.agents/specs/recurrent-multistate.md` named FOUR model families as consumers of `GdnStateCache::conv_state` / `ssm_state`, and the fourth reads neither field.** Fixed IN FLOW under W5c-1 of [#2031](https://github.com/mudler/vllm.cpp/issues/2031). The line listed `qwen3_5.cpp`, `kimi_linear_device.cpp`, `nemotron_h_device.cpp` and `gemma4_mm.cpp` as "Every existing consumer", to justify that widening `GdnStateCache` into an ordered `std::vector states` leaves them untouched. `gemma4_mm.cpp` has ZERO occurrences of `conv_state` and ZERO of `ssm_state`; its only two mentions of the type are an include comment (`:34`) and `std::vector no_gdn_state;` (`:221`), passed EMPTY. It is the file that proves Gemma-4 has no recurrent arm, cited as the file that proves the opposite — and `muse_glimmer_mm.cpp:340` and `qwen3_vl.cpp:621` carry that identical empty-vector shape, so the wrong fourth name was one of the three files demonstrating the negative. Measured at `ad6696fa3`, `GdnStateCache`/`conv_state`/`ssm_state` counts are 37/33/34 for `qwen3_5.cpp`, 6/9/14 for `nemotron_h_device.cpp`, 2/7/6 for `kimi_linear_device.cpp`, and 2/0/0 for each of the three non-consumers. The real count is THREE families: `qwen3_5`, `kimi_linear_device.cpp`, and `nemotron_h` (`nemotron_h_device.cpp` with `nemotron_h_forward.h`), each gathering and scattering through the named fields (`nemotron_h_device.cpp:1689-1690`, `kimi_linear_device.cpp:1774-1777`). THE NEAR-MISS THAT HIDES IT: grepping the FIELD name over-counts instead, because `glm5_next_kda.cpp:343-345` matches `conv_state` 13 times on `Glm5NextKdaCache::conv_state`, a `std::vector` KDA sequence state (`glm5_next_kda.h:314`) and not the `vt::Tensor` at `qwen3_5.h:111`, with zero occurrences of `GdnStateCache` — so the grep that under-counts is the one on the TYPE. WHAT IT COST: `f7710c1b4` ([#2131](https://github.com/mudler/vllm.cpp/issues/2131)) landed the line and W5c-1 copied the same four names into a shipped product comment (`src/vllm/model_executor/models/qwen4_exp_registry.cpp`) and into [`qwen4-exp-flash-next.md`](../specs/qwen4-exp-flash-next.md), both times as the justification for the deliberate `[gdn_conv, temporal, ple_conv, ngram]` state-order divergence. The CONCLUSION is unaffected at all three sites — moving the temporal state off slot 1 still silently re-points every consumer — but the blast radius written beside it was one family too wide and named a file whose behaviour is the opposite of the asserted one. No checker can see this class: `check-symbol-anchors` resolves SYMBOLS, and `GdnStateCache` genuinely appears in `gemma4_mm.cpp`, so symbol existence passes and a reader who greps the type finds the file and stops. Same defect class as [#2198](https://github.com/mudler/vllm.cpp/issues/2198), which W5c-1 also closes: a citation naming something that is not there, landing green because the thing it names exists somewhere nearby | bug | -| [#2178](https://github.com/mudler/vllm.cpp/issues/2178) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **No llama.cpp RELEASE defines `glm5next`, so register a scoped PR-pinned oracle — and the two candidate PRs turned out to be COMPETING implementations that disagree on the architecture string, not the text half and the vision half of one stack.** Registers [`llama-cpp-glm5next`](../oracles/llama-cpp-glm5next.md) at `ggml-org/llama.cpp` PR #27752, object `8a8d0bcc4d5fdf024c457526245bec4bc3a12adc`, on the `llama-cpp-qwen4exp` precedent. Re-measured 2026-08-28 in a fresh bare clone whose only remote is `ggml-org/llama.cpp`, from refs and objects and never from a working tree: `ls-remote` heads `8a8d0bcc...` (#27752) and `9370c82d...` (#27773) agree with `gh api .head.sha`; `fetch --depth 1` serves both; `merge-base --is-ancestor refs/heads/master` is **rc=1** for both against a `b10451` control at rc=0; `git grep -il 'glm5next\|glm5_next' b10451` is **rc=1** tree-wide against a `glm4_moe` control returning nine files, and the same grep at `master` `50f068fff` is rc=1 too; `conversion/glm5next.py` is 4714 B and `src/models/glm5next.cpp` 55716 B at the pin, against a `no-such-file.py` probe at rc=128. **#27752 registers `LLM_ARCH_GLM5NEXT -> "glm5next"` (`src/llama-arch.cpp:87`) and has no vision at all (`grep -il glm5 -- tools/` rc=1); #27773 registers `LLM_ARCH_GLM5_NEXT -> "glm5-next"` (`:152`) with its own text graph `src/models/glm5-next.cpp` plus `PROJECTOR_TYPE_GLM5V -> "glm5v"`.** The published `unsloth/GLM-5.3-Flash-GGUF` at revision `d425e572fb96` declares `general.architecture = glm5next` in its first shard's header, which is #27752's spelling and our own converter's, so pinning #27773 would give a denominator that refuses both artifacts by name — one file, not two. **O4 corrected** in [`glm5-next-flash.md`](../specs/glm5-next-flash.md): the RELEASE half holds, the "no llama.cpp oracle" half no longer does, and what stays owed is the floor itself plus a vision denominator. **W6's vision denominator is owed and #27773 would not discharge it even out of draft:** the staged `mmproj-BF16.gguf` declares `clip.projector_type = glm5next` and `grep -c '"glm5next"' -- tools/` is rc=1 at BOTH heads, so no revision of llama.cpp can load the published mmproj today. `gateable = no` with #2178 owing the measurement: nothing was built and nothing was run, and a build is not a run. The run half is REACHABLE for the first time — UD-Q2_K_XL (101.2535 GiB over four shards, summed on the `UD-Q2_K_XL/` prefix rather than a substring match, which also catches a 9,429,920-byte `Shard_Rewrite/` sibling that is not a shard) was staging to the NAS when this row was written; the oracle file carries the per-shard state at a named instant because a live count in an append-only row is a drift-lock. `mmproj-BF16.gguf` is complete at sha256 `513c9bfc55898998186543caefc01626fb28e378b92f391018e1c3dd6655b113` computed locally. **The contrast worth carrying:** the opposite case landed the same day in [#2194](https://github.com/mudler/vllm.cpp/issues/2194) — for `glm_moe_dsa` stock `b10451` ALREADY carries `LLM_ARCH_GLM_DSA -> "glm-dsa"` (`src/llama-arch.cpp:85`, case `:1051`, enumerator `src/llama-arch.h:90`, graph `src/models/glm-dsa.cpp`, converter `conversion/glm.py:274-276`), re-verified in this same clone at rc=0, so that row needed no scoped file at all. The test is one command, not a judgement: does the pinned RELEASE name the architecture? Records only: no product code, no pin advance on `llama-cpp`, no build, no GPU lease | feature | -| [#2218](https://github.com/mudler/vllm.cpp/issues/2218) | `MODEL-MM-QWEN4-EXP` | **The `hc_norm` gamma polarity disagrees between the loader and the device op, and a layer loop wiring them together scales by ~0.** `LoadNormBf16(..., unshift=true)` at `qwen4_exp_weights.cpp:264` stores the RAW HuggingFace gamma, centred on 0, by inverting the GGUF converter's baked `+1`. `vt::Qwen4ExpGatedResidual` documents the OPPOSITE convention — "hc_norm_w is vLLM's parameterization, i.e. ALREADY `1 + w_hf` … This op never adds 1" — so a layer loop that hands the loader's tensor straight to that op applies a near-zero scale, and the result reads as a corrupt checkpoint rather than as a wiring bug. The contradiction is visible AT THE LOAD SITE: the comment at `qwen4_exp_weights.cpp:258-263` argues FOR the fold, elementwise-corroborated on three published artifacts, immediately above the line that strips it. Nothing is broken today because `Qwen4ExpTextModel::Forward` does not exist; the moment the layer loop lands it must fold `hc_norm`, `norm_key`, `norm_query` and `norm_conv` through `vllm::qwen4_exp::HcNormWeightFromHf` first. NOT repaired in W5b-5, which hit the same shape and got it right by accident of scope: the QSA block's norms take the raw gamma under `RmsNormArgs::gemma = true`, which mutations M9/M10/M11 red. Owned by `MODEL-MM-QWEN4-EXP` and listed under `## Owed` in [`specs/qwen4-exp-flash-next.md`](../specs/qwen4-exp-flash-next.md); the layer-loop wave is where it gets fixed and gated. | bug | +| [#2203](https://github.com/mudler/vllm.cpp/issues/2203) | `MODEL-MM-QWEN4-EXP` | **`.agents/specs/recurrent-multistate.md` named FOUR model families as consumers of `GdnStateCache::conv_state` / `ssm_state`, and the fourth reads neither field.** Fixed IN FLOW under W5c-1 of [#2031](https://github.com/mudler/vllm.cpp/issues/2031). The line listed `qwen3_5.cpp`, `kimi_linear_device.cpp`, `nemotron_h_device.cpp` and `gemma4_mm.cpp` as "Every existing consumer", to justify that widening `GdnStateCache` into an ordered `std::vector states` leaves them untouched. `gemma4_mm.cpp` has ZERO occurrences of `conv_state` and ZERO of `ssm_state`; its only two mentions of the type are an include comment (`:34`) and `std::vector no_gdn_state;` (`:221`), passed EMPTY. It is the file that proves Gemma-4 has no recurrent arm, cited as the file that proves the opposite — and `muse_glimmer_mm.cpp:340` and `qwen3_vl.cpp:621` carry that identical empty-vector shape, so the wrong fourth name was one of the three files demonstrating the negative. Measured at `ad6696fa3`, `GdnStateCache`/`conv_state`/`ssm_state` counts are 37/33/34 for `qwen3_5.cpp`, 6/9/14 for `nemotron_h_device.cpp`, 2/7/6 for `kimi_linear_device.cpp`, and 2/0/0 for each of the three non-consumers. The real count is THREE families: `qwen3_5`, `kimi_linear_device.cpp`, and `nemotron_h` (`nemotron_h_device.cpp` with `nemotron_h_forward.h`), each gathering and scattering through the named fields (`nemotron_h_device.cpp:1689-1690`, `kimi_linear_device.cpp:1774-1777`). THE NEAR-MISS THAT HIDES IT: grepping the FIELD name over-counts instead, because `glm5_next_kda.cpp:343-345` matches `conv_state` 13 times on `Glm5NextKdaCache::conv_state`, a `std::vector` KDA sequence state (`glm5_next_kda.h:314`) and not the `vt::Tensor` at `qwen3_5.h:111`, with zero occurrences of `GdnStateCache` — so the grep that under-counts is the one on the TYPE. WHAT IT COST: `f7710c1b4` ([#2131](https://github.com/mudler/vllm.cpp/issues/2131)) landed the line and W5c-1 copied the same four names into a shipped product comment (`src/vllm/model_executor/models/qwen4_exp_registry.cpp`) and into [`qwen4-exp-flash-next.md`](specs/qwen4-exp-flash-next.md), both times as the justification for the deliberate `[gdn_conv, temporal, ple_conv, ngram]` state-order divergence. The CONCLUSION is unaffected at all three sites — moving the temporal state off slot 1 still silently re-points every consumer — but the blast radius written beside it was one family too wide and named a file whose behaviour is the opposite of the asserted one. No checker can see this class: `check-symbol-anchors` resolves SYMBOLS, and `GdnStateCache` genuinely appears in `gemma4_mm.cpp`, so symbol existence passes and a reader who greps the type finds the file and stops. Same defect class as [#2198](https://github.com/mudler/vllm.cpp/issues/2198), which W5c-1 also closes: a citation naming something that is not there, landing green because the thing it names exists somewhere nearby | bug | +| [#2178](https://github.com/mudler/vllm.cpp/issues/2178) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **No llama.cpp RELEASE defines `glm5next`, so register a scoped PR-pinned oracle — and the two candidate PRs turned out to be COMPETING implementations that disagree on the architecture string, not the text half and the vision half of one stack.** Registers [`llama-cpp-glm5next`](oracles/llama-cpp-glm5next.md) at `ggml-org/llama.cpp` PR #27752, object `8a8d0bcc4d5fdf024c457526245bec4bc3a12adc`, on the `llama-cpp-qwen4exp` precedent. Re-measured 2026-08-28 in a fresh bare clone whose only remote is `ggml-org/llama.cpp`, from refs and objects and never from a working tree: `ls-remote` heads `8a8d0bcc...` (#27752) and `9370c82d...` (#27773) agree with `gh api .head.sha`; `fetch --depth 1` serves both; `merge-base --is-ancestor refs/heads/master` is **rc=1** for both against a `b10451` control at rc=0; `git grep -il 'glm5next\|glm5_next' b10451` is **rc=1** tree-wide against a `glm4_moe` control returning nine files, and the same grep at `master` `50f068fff` is rc=1 too; `conversion/glm5next.py` is 4714 B and `src/models/glm5next.cpp` 55716 B at the pin, against a `no-such-file.py` probe at rc=128. **#27752 registers `LLM_ARCH_GLM5NEXT -> "glm5next"` (`src/llama-arch.cpp:87`) and has no vision at all (`grep -il glm5 -- tools/` rc=1); #27773 registers `LLM_ARCH_GLM5_NEXT -> "glm5-next"` (`:152`) with its own text graph `src/models/glm5-next.cpp` plus `PROJECTOR_TYPE_GLM5V -> "glm5v"`.** The published `unsloth/GLM-5.3-Flash-GGUF` at revision `d425e572fb96` declares `general.architecture = glm5next` in its first shard's header, which is #27752's spelling and our own converter's, so pinning #27773 would give a denominator that refuses both artifacts by name — one file, not two. **O4 corrected** in [`glm5-next-flash.md`](specs/glm5-next-flash.md): the RELEASE half holds, the "no llama.cpp oracle" half no longer does, and what stays owed is the floor itself plus a vision denominator. **W6's vision denominator is owed and #27773 would not discharge it even out of draft:** the staged `mmproj-BF16.gguf` declares `clip.projector_type = glm5next` and `grep -c '"glm5next"' -- tools/` is rc=1 at BOTH heads, so no revision of llama.cpp can load the published mmproj today. `gateable = no` with #2178 owing the measurement: nothing was built and nothing was run, and a build is not a run. The run half is REACHABLE for the first time — UD-Q2_K_XL (101.2535 GiB over four shards, summed on the `UD-Q2_K_XL/` prefix rather than a substring match, which also catches a 9,429,920-byte `Shard_Rewrite/` sibling that is not a shard) was staging to the NAS when this row was written; the oracle file carries the per-shard state at a named instant because a live count in an append-only row is a drift-lock. `mmproj-BF16.gguf` is complete at sha256 `513c9bfc55898998186543caefc01626fb28e378b92f391018e1c3dd6655b113` computed locally. **The contrast worth carrying:** the opposite case landed the same day in [#2194](https://github.com/mudler/vllm.cpp/issues/2194) — for `glm_moe_dsa` stock `b10451` ALREADY carries `LLM_ARCH_GLM_DSA -> "glm-dsa"` (`src/llama-arch.cpp:85`, case `:1051`, enumerator `src/llama-arch.h:90`, graph `src/models/glm-dsa.cpp`, converter `conversion/glm.py:274-276`), re-verified in this same clone at rc=0, so that row needed no scoped file at all. The test is one command, not a judgement: does the pinned RELEASE name the architecture? Records only: no product code, no pin advance on `llama-cpp`, no build, no GPU lease | feature | +| [#2218](https://github.com/mudler/vllm.cpp/issues/2218) | `MODEL-MM-QWEN4-EXP` | **The `hc_norm` gamma polarity disagrees between the loader and the device op, and a layer loop wiring them together scales by ~0.** `LoadNormBf16(..., unshift=true)` at `qwen4_exp_weights.cpp:264` stores the RAW HuggingFace gamma, centred on 0, by inverting the GGUF converter's baked `+1`. `vt::Qwen4ExpGatedResidual` documents the OPPOSITE convention — "hc_norm_w is vLLM's parameterization, i.e. ALREADY `1 + w_hf` … This op never adds 1" — so a layer loop that hands the loader's tensor straight to that op applies a near-zero scale, and the result reads as a corrupt checkpoint rather than as a wiring bug. The contradiction is visible AT THE LOAD SITE: the comment at `qwen4_exp_weights.cpp:258-263` argues FOR the fold, elementwise-corroborated on three published artifacts, immediately above the line that strips it. Nothing is broken today because `Qwen4ExpTextModel::Forward` does not exist; the moment the layer loop lands it must fold `hc_norm`, `norm_key`, `norm_query` and `norm_conv` through `vllm::qwen4_exp::HcNormWeightFromHf` first. NOT repaired in W5b-5, which hit the same shape and got it right by accident of scope: the QSA block's norms take the raw gamma under `RmsNormArgs::gemma = true`, which mutations M9/M10/M11 red. Owned by `MODEL-MM-QWEN4-EXP` and listed under `## Owed` in [`specs/qwen4-exp-flash-next.md`](specs/qwen4-exp-flash-next.md); the layer-loop wave is where it gets fixed and gated. | bug | | [#2213](https://github.com/mudler/vllm.cpp/issues/2213) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **NoPE MLA and the DSA k-pool indexer — the geometry every later wave waits on.** W3 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998). Two things, and each one fails quietly. (1) `MlaBlockDims::Validate` required every dimension `> 0` while `Glm5NextTextConfig.validate_architecture` REQUIRES `qk_rope_head_dim == 0` ("Expecting NoPE for the DSA attention layers"), so the two validators were exact complements over one field and no value satisfied both — O11, pinned executably in `test_glm5_next_scaffold.cpp` and now discharged: 0 is the ABSENT rotary, `head_size()` collapses to `kv_lora_rank` (512, not 576), and the block's rope branches become NOT TAKEN rather than zero-width work. Kimi-Linear is the near miss and is untouched: it keeps `qk_rope_head_dim = 64` and skips only the rotation. (2) `Glm5NextTextIndexer` scores LEARNED POOLED candidates, not raw tokens — `index_kpool` consecutive valid tokens compressed by a per-channel 4-way softmax with an intra-pool position embedding, `index_topk // index_kpool` pools selected, expanded back to raw indices, and the ragged tail appended raw and UNSCORED at width `index_topk + index_kpool - 1` = 2051. `deepseek_v4_dsa.cpp` has no pooling stage at all, so reusing it selects the wrong candidate set and yields plausible indices either way. `index_kpool` is **4** on the published artifact and 16 in the config class. Landed `src/vllm/model_executor/models/glm5_next_dsa.{h,cpp}` gated against goldens RUN out of transformers v5.16.1 at seq_len 21 vs index_topk 8 — STRICTLY past the threshold, because at or below it a top-k selects everything and the pooling is unobservable — asserting SET equality of the selected indices over 17 discriminating rows with a smallest margin of 2.58e-3. SACRED inertness proven by the six-arm DeepSeek byte-identity probe, base `150b37852` vs head, all six fingerprints identical | feature | | [#2223](https://github.com/mudler/vllm.cpp/issues/2223) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The 288 routed + 1 shared expert MoE, and the heterogeneous KV-cache spec — the first piece of this row that a production entry point REACHES.** W5 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998). Two deliverables. (1) `src/vllm/model_executor/models/glm5_next_moe.{h,cpp}` BINDS rather than reimplements: the router is `vt::MoeRouterTopK`'s grouped `noaux_tc` arm and the epilogue is `deepseek_v4::ClampedSwiGLU` at `alpha=1, beta=0`, which is `_apply_gate`'s "Simple swiglu instead of alpha" line for line. Gated at the PUBLISHED 288/top-8 against goldens RUN out of `Glm5NextTextTopkRouter.forward` at transformers `v5.16.1`, asserting SET equality of the selected experts and PRINTING the separation margin (smallest 1.84e-3 over four tokens) — top-k error is bimodal, so a tolerance passes a wrong selection whose values happen to be close. Five silent-failure axes are each a killed mutation: sigmoid vs softmax scoring, `routed_scaling_factor` dropped, `norm_topk_prob` dropped, the `e_score_correction_bias` dropped (the bias SELECTS, the unbiased score WEIGHTS), and the factor applied TWICE by also passing it to `vt::MoeCombine`. (2) `MakeGlm5NextKVCache` replaces a refusal with THREE published groups — an `MLAAttentionSpec` at **512** for the 11 DSA layers (NOT the 576 every DeepSeek variant and Kimi-Linear publish: `qk_rope_head_dim` is ZERO here and upstream requires it), ONE uniform `MambaSpec` for the 34 KDA layers whose conv state is `conv_kernel_dim` = **4** columns wide and NOT `K-1` (`cache_utils.py:1015-1024` allocates it that wide and `causal_conv1d_update` reads `state_len` back off it; `kimi_linear_registry.cpp:157` publishes `K-1` for ITS model and copying that across hands the runner a cache one column short), and an `MLAAttentionSpec` at **257** = `2*index_head_dim + 1` for the indexer side cache at `compress_ratio` **1**, because the k-pool compresses at READ time and not in the store — the opposite of `MODEL-MM-QWEN4-EXP`'s QSA side cache. **REACHED**: the cases enter through `ModelRegistry::Resolve` and the `make_kv_cache` factory hook; unwiring `.make_kv_cache` REDS the gate, and DELETING the row is a `-Werror=unused-function` build error, so the toolchain itself proves the factory is the only reference. Measured on the way, and SUPERSEDED before this wave landed: the production loader run against the staged `unsloth/GLM-5.3-Flash-GGUF` rev `d425e572f` UD-Q2_K_XL arm opened the file, resolved `glm5next`, walked the 4-way split and stopped on `blk.3.ffn_gate_exps.weight has unknown ggml type id 17` (IQ2_XS). That was true when W5 measured it on 2026-08-29 and is not true now: [#2245](https://github.com/mudler/vllm.cpp/issues/2245) landed the IQ2_XS and IQ4_XS decoders and W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)) resolves all 1383 backbone tensors of that artifact, so the reading is kept as the measurement it was rather than as a live claim. The decoder layer, the DSA attention block and the assembled text forward are NOT in this wave and are carried as O23 | feature | | [#2230](https://github.com/mudler/vllm.cpp/issues/2230) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **Three refusal messages named LANDED waves as owing, and one denied an artifact that exists — and the gate was PINNING all three.** Fixed IN FLOW under W5 of [#2223](https://github.com/mudler/vllm.cpp/issues/2223). (1) The forward refusal read "W3 the NoPE MLA block -- `MlaBlockDims::Validate` still refuses `qk_rope_head_dim == 0`", which W3 (#2213, `e511a614b`) made false by relaxing exactly that validator; it named W2's sigmoid forget gate and W4's unweighted mHC collapse as owed too, both landed (`199c44578`, `6c715de00`). W1 wrote the message and no later wave touched the file — `git log --oneline -- src/vllm/model_executor/models/glm5_next_registry.cpp` ends at W1's `47a2b35a5`. (2) The GGUF loader refusal read "NO `.gguf` of this model exists anywhere ... (O7)"; `unsloth/GLM-5.3-Flash-GGUF` rev `d425e572f` is published and four arms are staged. (3) The KV-cache refusal said the KDA layers carry "three separate conv states"; they carry ONE — the checkpoint's `self_attn.{q,k,v}_conv1d` concatenate into one grouped depthwise conv (`modeling_glm5_next.py:620-628`, `glm5_next_kda.h` "THREE LAYOUT FACTS"), and a spec written from that sentence would triple the group. THE MECHANISM: `test_glm5_next_scaffold.cpp` asserted all three sentences, so the gate passed *because* nothing had corrected them — a refusal message is this row's only user-visible surface and its assertions were pinning stale text rather than checking it. The repair adds the NEGATIVES (`MlaBlockDims::Validate still refuses` and `NO `.gguf` of this model exists` must NOT appear) so a revision that reintroduces either reds | bug | | [#2234](https://github.com/mudler/vllm.cpp/issues/2234) | `SPEC-DFLASH2` | **The batched-lane spec's `## Now` told a reader that L2 must NOT be merged, and L2 had been on `main` since `150b37852`.** `scripts/now.py` renders a row's live position from `## Now`, so the derived surface reported a landed change (#2212) as an unmergeable branch — the same defect class as [#2199](https://github.com/mudler/vllm.cpp/issues/2199), where a section written before a wave landed was never reconciled by the landing. Record-only repair, no product code. Three further claims had drifted and are marked DISCHARGED in place rather than deleted, so a later reader can tell "done" from "never written": the seam policy item landed as `c9b2049bc` (#2207), which is what makes a quantized gate-up arm reachable for the draft at all and so is a precondition of [#2224](https://github.com/mudler/vllm.cpp/issues/2224); O3 was already closed in `dflash2-batch-propose.md:348`; and the stale-anchor bullet cited the `P == 1` gate as `:1614` when it is `:1716`, so the CORRECTION had drifted twice as far as the `:1577` it was written to fix, which is the argument for `.agents/porting.md`'s name-the-symbol rule stated twice over. `## Now` now records L2's measured **-11% on `fwd`** (35.19 -> 31.3 ms, terminal control matching to 1.1%) and states **L3, the batched capture lane, as the row's next gate**: at `P > 1` the draft forward is not capture-targeted, so at c=8 the term that is 76% of the draft phase runs EAGER, while vLLM replays a FULL draft graph at every batch size and pads to `max_num_reqs` with `PAD_SLOT_ID` — verified at the parity pin `5559679229`, `spec_decode/dflash/speculator.py:456-458` (`run_fullgraph`) and `:589` ("Pad per-request buffers to max_num_reqs for CUDA graph safety"). A porting gap under "mirror vLLM", not a new design. It also records that the binaries carry no tree identity — `vllm_version()` returns `0.0.3+cuda` for every commit because `VLLM_CPP_BUILD_VERSION` defaults to `PROJECT_VERSION` — so L2's build is identified by its KERNEL SIGNATURE instead (`DFlashAttnMmaKernel` mangling to `...fbll`, 13 params carrying `tiles_per_req`, against `...fbl` in the pre-L2 `build23`), which proves the feature is compiled in rather than that a directory was named after a SHA | bug | | [#2240](https://github.com/mudler/vllm.cpp/issues/2240) | `QUANT-GGUF-IQ2_XS` | **IQ2_XS (17) and IQ4_XS (23) — the last two GGUF dequantizers the staged GLM-5.3-Flash artifact needed, and the two the loader stopped dead on.** "UD-Q2_K_XL" names a target average, not a format: of that artifact's 1412 tensors only TWO are Q2_K, while 82 are IQ2_XS (the `ffn_gate_exps`/`ffn_up_exps` routed experts) and 3 are IQ4_XS, so `LoadedEngine::FromModelDir` refused at `blk.3.ffn_gate_exps.weight has unknown ggml type id 17` before any dequant code ran — the reader had no block stride for 17, and the switch had no decoder for either. Both ported 1:1 from llama.cpp `b10451` (`ggml/src/ggml-quants.c:2516` `dequantize_row_iq2_xs`, `:2743` `dequantize_row_iq4_xs`, `ggml/src/ggml-common.h:627` `iq2xs_grid`) and gated BYTE-FOR-BYTE against the oracle's own decoders over REAL bytes read out of the two tensors that failed. IQ2_XS is the middle member of a family of three same-shaped codebooks — 256 / 512 / 1024 entries — where reaching for the wrong table still runs and still produces plausible magnitudes, so the 512-entry grid carries an FNV-1a seal as well. IQ4_XS reuses `kValuesIq4nl` unchanged; its delta is the super-block scale layout, a 6-bit `ls` spliced from a `scales_l` nibble and a `scales_h` bit pair and then biased by -32. Also carries the record correction the issue asked for: `.agents/specs/glm5-next-flash.md` O5/O8 are about the converter's WRITE side and were read as meaning the i-quant lane was absent entirely. Owning row `QUANT-GGUF-IQ2_XS` (and `QUANT-GGUF-IQ4_XS`); found by W5 of [#1998](https://github.com/mudler/vllm.cpp/issues/1998) via [#2223](https://github.com/mudler/vllm.cpp/issues/2223) | feature | -| [#2243](https://github.com/mudler/vllm.cpp/issues/2243) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **`glm5next.attention.head_count_kv` is a per-layer `array[i32]` in the published artifact and `Glm5NextHfConfigFromGguf` reads it as a scalar.** Found while landing [#2240](https://github.com/mudler/vllm.cpp/issues/2240): with IQ2_XS and IQ4_XS decoded, the production loader gets past the type-17 refusal, opens all four shards, sizes all 1412 tensors, and stops instead at `glm5_next gguf: key glm5next.attention.head_count_kv is not an integer`. The artifact stores the layer schedule there — length 46, `0` on the 35 KDA layers and `1` on the 11 DSA/MLA layers — and `swiglu_clamp_exp`/`swiglu_clamp_shexp` are per-layer `array[f32]` of the same length directly behind it. Filed rather than fixed in that flow because it belongs to this row's config/loader wave and not to a dequant change; listed under `## Owed` as O18 in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) | bug | +| [#2243](https://github.com/mudler/vllm.cpp/issues/2243) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **`glm5next.attention.head_count_kv` is a per-layer `array[i32]` in the published artifact and `Glm5NextHfConfigFromGguf` reads it as a scalar.** Found while landing [#2240](https://github.com/mudler/vllm.cpp/issues/2240): with IQ2_XS and IQ4_XS decoded, the production loader gets past the type-17 refusal, opens all four shards, sizes all 1412 tensors, and stops instead at `glm5_next gguf: key glm5next.attention.head_count_kv is not an integer`. The artifact stores the layer schedule there — length 46, `0` on the 35 KDA layers and `1` on the 11 DSA/MLA layers — and `swiglu_clamp_exp`/`swiglu_clamp_shexp` are per-layer `array[f32]` of the same length directly behind it. Filed rather than fixed in that flow because it belongs to this row's config/loader wave and not to a dequant change; listed under `## Owed` as O18 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) | bug | | [#2201](https://github.com/mudler/vllm.cpp/issues/2201) | `BACKEND-TENSTORRENT-QWEN35` | **W3, the GDN row's reviewer leftovers: the d2h counter misses two download paths, and `EnsureGdnCacheDevice`'s fast path accepts a conv-transposed host pointer.** `GdnStateD2hBytes()` (`tenstorrent_ops.cpp:4164`) counts `:5039`/`:5109` but not the `EnsureGdnCacheDevice` (`:4216`) slow-path download nor the `CommitConvTransposed` (`:4563`) untracked-buffer fallback, so `state_d2h_bytes` is a lower bound and counter-asserting legs cannot see those paths; and the fast path keys on the host pointer alone without the `conv_transposed` role check, so cross-role pointer reuse would serve a wrong-geometry cached tensor (not live today — `qwen3_5.cpp` uses distinct buffers). Test-first fix in `tests/vt/test_tenstorrent_backend.cpp`: red per leftover, then the two missing `fetch_add`s and a role-mismatch refusal that names it; sacred 16/16 goldens byte-identical | bug | | [#2214](https://github.com/mudler/vllm.cpp/issues/2214) | `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` | **GLM-5.3 is 97.49% routed experts, so the hardware blocker was computed in the wrong FRAME and the row moves `BLOCKED` -> `SPIKE` on a committed port plan.** [#2194](https://github.com/mudler/vllm.cpp/issues/2194) sized RESIDENT weights — 753,329,940,480 parameters, 703.74 GiB of fp8, 1403.2 GiB bf16 against 119.631 GiB on `dgx:gpu0`, needing 1.36 bpw — which is the right frame for a dense model and the wrong one for a 256-expert MoE. Recomputed here from `zai-org/GLM-5.3`'s own `config.json` and checked against the checkpoint's `model.safetensors.index.json` (`metadata.total_size = 755,617,140,416` over 118,629 tensors): routed experts are **734,439,407,616 of 753,328,717,824 = 97.49% streamable**, resident is **18,889,310,208 = 2.51%**, and the total reproduces the measured `safetensors.total` to **-1,222,656 params (-0.00016%)** — tighter than #2214's own 745.8B model, which omitted the MTP block's 256 experts and put resident 2.1B too high. **The decisive number is measured from the artifact, not derived.** A full GGUF header census of `unsloth/GLM-5.3-GGUF` `UD-IQ1_S` at revision `346b3591c7f28d1a23716f97a065ecf12ec14771` — 6 shards, 1809 tensors matching the file's own `split.tensors.count`, read by HTTP RANGE request at a cost of ~9.6 MB with nothing downloaded — splits it into **228 `*_exps.weight` towers = 187.312 GiB streamable and 1581 resident tensors = 14.511 GiB**, and the resident class is ~14.5 GiB in every published UD arm because the recipe keeps every non-expert tensor at Q4_K or better. One decode step at `c = 1` touches `75 x 3 x 8 = 1800` distinct slices = **11.21 GiB** of uniform slots, so resident + a 4096-slot cache is **40.01 GiB**. **A `UD-*` name is a target average and not a format:** `UD-IQ1_S` is 106 IQ1_S + 71 IQ3_XXS + 44 IQ2_XXS + 4 IQ4_XS + 3 K-quant expert tensors, and `UD-Q2_K_XL` holds exactly TWO Q2_K tensors out of 1809. **The decoder/`vec_dot` verdict was re-measured against the merged tree, because `origin/main` moved under this branch and inverted it.** At the base `60a6dd97b`, neither `IQ4_XS` (23) nor `IQ2_XS` (17) had a vt block dtype or a decoder and both were a hard refusal. `94de63ff5` ([#2245](https://github.com/mudler/vllm.cpp/issues/2245)), landed 2026-08-29 for the sibling Flash row, ported BOTH DEQUANTIZERS and NEITHER keep-quant `vec_dot`. `gguf_keep_quant.cpp::KeepQuantDType` returns false without `vt::cpu::HasQuantDotKernel`, so **a type with a decoder and no `vec_dot` EXPANDS to bf16 at load**, and `gguf_device_fit.cpp:85-100` is all-or-nothing across a model's `*_exps` tensors, so four offending tensors out of 228 drop the whole arm out of the expert-streaming lane. One `*_exps` tower is 3,221,225,472 elements = 6.000 GiB at bf16, so `UD-IQ1_S`'s four IQ4_XS towers go **6.375 -> 24.000 GiB** and `UD-Q2_K_XL`'s 148 IQ2_XS towers go **128.344 -> 888.000 GiB**, while the uniform slot grows from 6.375 MiB to a 24.00 MiB bf16 slice. **So the row is blocked on exactly one kernel and it is a `vec_dot`, not a decoder: `VecDotIQ4_XSQ8_K`**, four tensors `blk.{8,75,76,77}.ffn_down_exps.weight`, which also unlocks `UD-IQ2_M`; `QUANT-GGUF-IQ4_XS` already owns it. `IQ1_M` (29) still has no reader traits, so `UD-IQ1_M` refuses at file OPEN. The general defect, named but not repaired here: landing a decoder without its `vec_dot` converts a loud refusal into a silent 3.4x memory multiplier no token gate can see. **One premise of #2194 is corrected: the pinned vLLM class CAN load this checkpoint.** It never reads `indexer_types` — zero occurrences tree-wide at `555967922` — it DERIVES the schedule at `deepseek_v2.py:1092-1103` from `index_topk_freq`/`index_skip_topk_offset`, giving 21 full trunk indexers plus the MTP layer, and drops surplus checkpoint indexer weights at `:1566-1582`. The config's `indexer_types`, that derivation, and llama.cpp `b10451`'s hardcoded `GLM_5_2_DEFAULT_INDEXER_TYPES` (`src/models/glm-dsa.cpp:6-27`) agree bit for bit over all 78 layers, and the checkpoint ships `self_attn.indexer.*` on exactly 22 of 79 blocks. `modules_to_not_convert`'s `self_attn.indexers_proj` matches no shipped tensor and no upstream module at the pin, so it is a config-level shorthand and a loader must not mirror it. **The delta is smaller than the spike's own DSA verdict implies**, because a device-native DSA lightning indexer now lives in the shared MLA block (`mla_attention.cpp:598-745`, CPU+CUDA `DsaIndexerLogits`/`DsaTopkSelect`, reached in production by `Dots3NoteForCausalLM`) and this model's MLA geometry already validates and dispatches to native kernels. Net-new: the indexer KV side cache ([#1925](https://github.com/mudler/vllm.cpp/issues/1925)), sparse prefill, the heterogeneous per-layer schedule with `skip_topk` selection reuse, the fp32 router GEMM, a `"glm-dsa"` `kGgufArchArms` row, and lifting the expert-streaming seam out of `qwen3_5.cpp`, which is the only model TU that constructs it. **The gate is the honest cost: NO end-to-end token gate against vLLM is reachable on this fleet**, because vLLM implements the architecture and cannot run 703.74 GiB on a 119.631 GiB unified device whose host RAM is the same pool; what is reachable is module parity against the pin on CPU, a headers-only structural loader gate, a streamed-vs-resident identical-logits gate needing no oracle, and an llama.cpp `b10451` floor on the identical artifact, with speed an open gap by construction. Eight waves planned, none landed. Records and spec only: no product code, no pin advance, no build, no GPU lease, no download | feature | -| [#2247](https://github.com/mudler/vllm.cpp/issues/2247) | `QUANT-GGUF-IQ2_XS` | **Keep-quant `vec_dot` for IQ2_XS and IQ4_XS: 325.58 GiB, and the difference between the staged GLM-5.3-Flash artifact fitting `dgx:gpu0` and overflowing it 3.6x.** [#2245](https://github.com/mudler/vllm.cpp/pull/2245) gave both types a row DECODER, which is what moved the loader past `unknown ggml type id 17`. A decode-only type has no `vec_dot`, so `HasQuantDotKernel` is false and every GEMM weight of that type expands to bf16 at load. Measured from the artifact's own headers, all four shards and all 1412 tensors: **101.24 GiB on disk, 597.46 GiB as bf16**, an expansion of 5.9x, of which IQ2_XS alone is 53.33 -> 369.00 GiB and IQ4_XS 3.59 -> 13.50 GiB. Resident TODAY **426.72 GiB** against the ~119.63 GiB the box has, so it does not fit; with these two kernels **101.14 GiB**, which fits with 18.49 GiB of headroom. Every other encoding in the file already keeps its quantization, IQ3_XXS (`VecDotIQ3_XXSQ8_K`) included, so these two are the entire gap. Two rows in `src/vt/cpu/cpu_quant_dot.cpp` beside the fifteen already there, ported from the pinned llama.cpp `b10451` and gated BYTE-FOR-BYTE against the oracle's own kernel on real artifact bytes, because a `vec_dot` defect shows up as numeric drift and not as a crash. Owning rows `QUANT-GGUF-IQ2_XS` and `QUANT-GGUF-IQ4_XS` in [`quantization-matrix.md`](../quantization-matrix.md), both carrying it as `C` = `-`; also recorded as O18 in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) | feature | +| [#2247](https://github.com/mudler/vllm.cpp/issues/2247) | `QUANT-GGUF-IQ2_XS` | **Keep-quant `vec_dot` for IQ2_XS and IQ4_XS: 325.58 GiB, and the difference between the staged GLM-5.3-Flash artifact fitting `dgx:gpu0` and overflowing it 3.6x.** [#2245](https://github.com/mudler/vllm.cpp/pull/2245) gave both types a row DECODER, which is what moved the loader past `unknown ggml type id 17`. A decode-only type has no `vec_dot`, so `HasQuantDotKernel` is false and every GEMM weight of that type expands to bf16 at load. Measured from the artifact's own headers, all four shards and all 1412 tensors: **101.24 GiB on disk, 597.46 GiB as bf16**, an expansion of 5.9x, of which IQ2_XS alone is 53.33 -> 369.00 GiB and IQ4_XS 3.59 -> 13.50 GiB. Resident TODAY **426.72 GiB** against the ~119.63 GiB the box has, so it does not fit; with these two kernels **101.14 GiB**, which fits with 18.49 GiB of headroom. Every other encoding in the file already keeps its quantization, IQ3_XXS (`VecDotIQ3_XXSQ8_K`) included, so these two are the entire gap. Two rows in `src/vt/cpu/cpu_quant_dot.cpp` beside the fifteen already there, ported from the pinned llama.cpp `b10451` and gated BYTE-FOR-BYTE against the oracle's own kernel on real artifact bytes, because a `vec_dot` defect shows up as numeric drift and not as a crash. Owning rows `QUANT-GGUF-IQ2_XS` and `QUANT-GGUF-IQ4_XS` in [`quantization-matrix.md`](quantization-matrix.md), both carrying it as `C` = `-`; also recorded as O18 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) | feature | | [#2252](https://github.com/mudler/vllm.cpp/issues/2252) | `SPEC-DFLASH2` | **The draft block's paged attention set neither `query_start_loc_host` nor `max_seq_len`, so it took a `cudaStreamSynchronize` path inside the CUDA graph capture it runs in, and the committed speed gate could not produce a number at all.** Found by running `scripts/dflash2-speed-gate.sh` on `main` at `e65678109` in an `rc` lease on `dgx:gpu0`: our arm exited 1 with `[vt cuda: paged prefill qsl sync: operation not permitted when stream is capturing]`, `GATE_RC=2`, `RESULT_PRESENT=no`. `DflashBlockPagedAttention` (`qwen3_dflash_internal.h`) set `scale`, `causal`, `window_size` and `uniform_spec_query_len` but left BOTH host-metadata fields at their defaults, and `include/vt/ops.h` states the cost of each in the same words — `query_start_loc_host` nullptr "=> the launcher falls back to the D2H+sync" (`:1546`), `max_seq_len` 0 "=> that launcher falls back to the D2H+sync" (`:1555`). That fallback is `BuildPrefillTiles`' legacy arm (`cuda_paged_attn.cu:2267-2272`), a D2H copy plus `cudaStreamSynchronize`, shared by all six prefill launchers (`:2301`, `:2333`, `:2371`, `:2411`, `:2452`, `:2495`), so which one ran does not change the outcome — **and the draft block is the ONE lane this tree captures** (`P == 1`, `qwen3_dflash.cpp:1716`, replaying `st.g_graph`). REGRESSION with a witness: `9aea9efec` (W11, [#1890](https://github.com/mudler/vllm.cpp/issues/1890)) added the call and is NOT an ancestor of `04ed7b984`, where the same gate's our-arm demonstrably RAN — the O28B evidence carries `our-arm.json` and `clock-ours.json` and refused only on a clock window ("idle for 2630 of 3222 SM-clock samples, 18.37% busy, below the 50% floor"), which is reachable only by executing and being sampled. WHY NOTHING CAUGHT IT: both fields are OPTIONAL by type, so omitting them compiles and is merely SLOW wherever nothing is capturing; the CPU backend ignores them entirely, so the file's own byte-identical equivalence suite stayed green through the whole defect (**measured**: under the mutation 13 of 15 cases still passed); and the failure needs CUDA AND capture AND the prefill ladder together, which nothing in CI reaches ([#1972](https://github.com/mudler/vllm.cpp/issues/1972)). Fixed by setting both from values the function already derives — the single-request host qsl `[0, tq)` and the EXTENDED bound `ctx_len + tq`, which `DflashBlockPagedInputsOf` computes and the call already refuses on if it disagrees; an upper bound is explicitly safe because it only sizes grids (`ops.h:1551-1553`). Gated by extracting the args build into a pure `DflashBlockPagedArgsOf`, on the `DflashBlockPagedInputsOf` / `DflashBlockPagedMaskOf` precedent this file already sets, because asserting the host-meta helper ALONE would pass while production forgot to use it — which is exactly how this shipped. Red-first proven by mutation: both new cases fail on `REQUIRE(pa.query_start_loc_host != nullptr)` with the fields defaulted, 48/48 assertions green with them set, and all 29 `dflash` test binaries pass | bug | | [#2250](https://github.com/mudler/vllm.cpp/issues/2250) | `MODEL-TEXT-laguna-laguna-for-causal-lm` | Lever #2: measure what bounds `QuantDotGemmGroupedKernel` (Q4_K/Q5_K, 62.1% of Laguna decode GPU) BEFORE tuning it. W11 labelled it "BW-tuning", but the sibling `QuantDotGemmQ8_0Kernel` was measured LATENCY- and LSU-pipe-bound with five structural levers refuted and a recorded floor, so the label is not evidence. W1 is `ncu` counters on the decode path only | perf | -| [#2177](https://github.com/mudler/vllm.cpp/issues/2177) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The GGUF config builder never read `glm5next.attention.head_count_kv`, so the published artifact's layer schedule was refused rather than read — and the `idx % 4 != 3` fallback that would have replaced it is right on this checkpoint by coincidence.** Fixed together with [#2243](https://github.com/mudler/vllm.cpp/issues/2243), because they are two faces of one defect: `Glm5NextHfConfigFromGguf` read that key as a scalar (`OptInt` -> `KvInt`, whose `default:` arm threw `key glm5next.attention.head_count_kv is not an integer`) and then REQUIRED a `glm5next.layer_types` string array only `scripts/convert-glm5-next-gguf.py` writes. The published `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL artifact carries neither shape: measured 2026-08-29 out of shard 1's own 72-key KV block, `attention.head_count_kv` is `array[i32]` n=46, `swiglu_clamp_exp` and `swiglu_clamp_shexp` are `array[f32]` n=46, and there is no `layer_types` key at all. The reader now accepts llama.cpp's scalar-or-array spelling of that key (`b10451:src/llama-model.cpp:1177` reads it through `get_key_or_arr(..., n_layer, false)`) and DERIVES the schedule from the values with llama.cpp's own predicate, `is_recr_impl[i] = hparams.n_head_kv(i) == 0` (`b10451:src/models/kimi-linear.cpp:18`, "KDA layers are recurrent"). **The array is 34 zeros and 12 ones, not 35 and 11**: `block_count` is 46 because it counts the MTP block (`nextn_predict_layers = 1`) on top of 45 model layers, the ones sit at 3, 7, ..., 43 AND at 45, and `45 % 4 == 1` — so a consumer that re-derives the stride selects eleven MLA-shaped blocks where the file states twelve and reports nothing. That is why the non-stride fixture, not the published one, is the case that proves the values are read. The array is a SCHEDULE and not a KV-head count: its non-zero entries are `1`, the single latent KV head MLA has, while upstream requires `num_attention_heads == num_key_value_heads`, so the array form leaves `num_key_value_heads` at upstream's `None -> num_attention_heads` default. When both spellings are present they are cross-checked on the layer KIND and a clash refuses by name; a per-layer array whose length is not `block_count` refuses by name with the shape found; a non-uniform clamp array refuses, because upstream has ONE `swiglu_limit`; and a file stating neither schedule still refuses, naming both keys. Listed as O18 in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) | bug | -| [#2268](https://github.com/mudler/vllm.cpp/issues/2268) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The published GGUF spells MLA geometry by llama.cpp's `attention.key_length` convention and carries no `attention.linear_head_count`, so the loader stops there next.** Measured 2026-08-29 on one tree and one binary, with the [#2243](https://github.com/mudler/vllm.cpp/issues/2243) / [#2177](https://github.com/mudler/vllm.cpp/issues/2177) array fix reverted and restored, driving `LoadedEngine::FromModelDir` at `/mnt/nas_share/rc/ckpt/GLM-5.3-Flash-UD-Q2_K_XL/` on `device = kCPU`, headers only: without the fix it stops at `key glm5next.attention.head_count_kv is not an integer`, with it at `attention.key_length_mla - attention.key_length is -256 but rope.dimension_count is 0` (`glm5_next_weights.cpp:435`). The file is not malformed. llama.cpp writes `key_length = kv_lora_rank + qk_rope_head_dim` and `key_length_mla = qk_nope_head_dim + qk_rope_head_dim` (`b10451:conversion/deepseek.py:345-348`), which for this model gives the artifact's own 512 / 512 / 256 / 256; `scripts/convert-glm5-next-gguf.py` writes `key_length = qk_nope_head_dim`, a different quantity under the same name. `glm5next.attention.linear_head_count`, a `ReqInt` in the builder, is in none of the file's 72 keys and llama.cpp spells it nowhere. Filed rather than fixed in that flow because it moves the WRITE side: putting the reader on llama.cpp's meaning without moving the converter would refuse our own output, and which spelling this project writes is a row-and-spec decision. O7 records that our converter has never been run, so no artifact of ours is invalidated by the move. Recorded under `## Owed` as O18 in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) | bug | -| [#2220](https://github.com/mudler/vllm.cpp/issues/2220) | `LTX25-ORACLE-ABSOLUTE` | **A CUDA toolkit staged off CIFS compiles but cannot be LINKED against, and the precondition that should catch it checks the one link that works.** `/workspace` is CIFS and stores no symlink, so a staged toolkit carries only `libcudart.so.13.3.29` and `libcublasLt.so.13.6.0.2`. `ltx25-oracle-absolute-render.sh` rebuilt the links with `b=${f%%.so.*}; ln -sf "$f" "$b.so"; ln -sf "$f" "$b.so.${f#*.so.}"` -- and `${f#*.so.}` strips the SHORTEST prefix, so it expands to `13.3.29` rather than `13`. The second `ln` therefore links `libcudart.so.13.3.29` to ITSELF and **`libcudart.so.13`, the SONAME, is never created**. That is the name `ld` resolves versioned undefined symbols against, so CMake reports `Found CUDAToolkit`, every CUDA TU compiles, and the job dies 21 minutes later with 38 `undefined reference to ...@libcudart.so.13` / `@libcublasLt.so.13` and `ninja: build stopped`. `need_ok` tested `[ -f .../libcublasLt.so ]`, which is exactly the link the loop DID create, so it passed on an unlinkable toolkit -- a precondition that cannot fail. LATENT, not new: the staging branch is a FALLBACK, and every earlier lease found `/usr/local/cuda` 13.0.88 and never took it; `dgx:gpu0` went `unhealthy ... worker_lost` for 3h20m on 2026-08-28 and returned without a toolkit, which exercised it for the first time. The two runs A/B in their own configure logs: `20260827T220845Z` `/usr/local/cuda` 13.0.88 built in 1192 s, `20260828T224529Z` `/root/cudatk` 13.3.73 failed at link. Fixed in flow: take the MAJOR (`v=${f#*.so.}; ${v%%.*}`), prefer `ldconfig -n` which reads each object's own `DT_SONAME`, and assert `.so` resolves AND `.so.` exists for both libraries BEFORE the build. Red-before/green-after on a replica of the CIFS layout: old loop creates no `.so.13`, new logic creates both, and the guard FAILS on the old layout, PASSES on the new, and FAILS on the real NAS source. `rc` job `1ad519b1-4e75-41d7-9386-9932076390f1`, exit 34. Also recorded in [`environment.md`](../environment.md) as a lease-environment fact, because it will bite the next row | bug | -| [#2186](https://github.com/mudler/vllm.cpp/issues/2186) | `MODEL-DSV4-EXL3` | **The DeepSeek-V4-Flash EXL3 artifact refused at 108.59 GiB because W1c materialized the carried tower to f32 from one- and two-byte storage, and 26.64 GiB of the total was that widening.** FIXED by W1d: the nine FP8-sourced fields of `DeepseekV4LayerHostWeights` (`wq_a`, `wq_b`, `wkv`, `wo_a`, `wo_b`, `idx_wq`, `shared_w1/w2/w3`) are held at `HostBf16` instead of `std::vector`, taking that half from 21.82 GiB to 10.91 GiB and the artifact's projected residency from 108.59 GiB to ~97.7 GiB against 119.63 GiB physical. **The narrowing is exactly lossless and the gate says so rather than accepting a tolerance**: E4M3 carries four significand bits and E8M0 is a pure power of two, so every value of this tower is exactly representable in bf16's eight, and `narrowing_lost_a_bit == 0` is asserted per element alongside the existing value equality. `Dot`, `MatVec`, `Gemm`, `GroupedOutputLora` and the CUDA `GroupedOLoraKernel` widen each weight AS THEY READ IT -- half the bytes moved, f32 accumulators, reduction order unchanged -- so nothing is materialized back to f32 on any path. **The residency accounting is now gated for the first time**: every other residency case in `test_deepseek_v4_exl3_loader.cpp` compares `DeepseekV4HostResidentBytes` against itself and stays green for any self-consistent formula, so a new case rebuilds the total from each loaded field's own `value_type`; hardcoding `sizeof(float)` back into the accounting takes exactly that one case red (MUTATION-PROVEN 2026-08-29). **Does NOT claim the artifact runs**: the DSA composition is still unported and the forward still refuses by name on the 21 `compress_ratio == 4` layers ([#1961](https://github.com/mudler/vllm.cpp/issues/1961), [#1970](https://github.com/mudler/vllm.cpp/issues/1970), [#1976](https://github.com/mudler/vllm.cpp/issues/1976)), and the ~97.7 GiB figure is a projection from the measured split, not a load that has been observed to complete. The remaining ~2.6 GiB of widening is the BF16-sourced norms/embeddings/router, left at f32 and still owed. Spec [model-dsv4-exl3.md](../specs/model-dsv4-exl3.md) `## W1d design` | bug | +| [#2177](https://github.com/mudler/vllm.cpp/issues/2177) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The GGUF config builder never read `glm5next.attention.head_count_kv`, so the published artifact's layer schedule was refused rather than read — and the `idx % 4 != 3` fallback that would have replaced it is right on this checkpoint by coincidence.** Fixed together with [#2243](https://github.com/mudler/vllm.cpp/issues/2243), because they are two faces of one defect: `Glm5NextHfConfigFromGguf` read that key as a scalar (`OptInt` -> `KvInt`, whose `default:` arm threw `key glm5next.attention.head_count_kv is not an integer`) and then REQUIRED a `glm5next.layer_types` string array only `scripts/convert-glm5-next-gguf.py` writes. The published `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL artifact carries neither shape: measured 2026-08-29 out of shard 1's own 72-key KV block, `attention.head_count_kv` is `array[i32]` n=46, `swiglu_clamp_exp` and `swiglu_clamp_shexp` are `array[f32]` n=46, and there is no `layer_types` key at all. The reader now accepts llama.cpp's scalar-or-array spelling of that key (`b10451:src/llama-model.cpp:1177` reads it through `get_key_or_arr(..., n_layer, false)`) and DERIVES the schedule from the values with llama.cpp's own predicate, `is_recr_impl[i] = hparams.n_head_kv(i) == 0` (`b10451:src/models/kimi-linear.cpp:18`, "KDA layers are recurrent"). **The array is 34 zeros and 12 ones, not 35 and 11**: `block_count` is 46 because it counts the MTP block (`nextn_predict_layers = 1`) on top of 45 model layers, the ones sit at 3, 7, ..., 43 AND at 45, and `45 % 4 == 1` — so a consumer that re-derives the stride selects eleven MLA-shaped blocks where the file states twelve and reports nothing. That is why the non-stride fixture, not the published one, is the case that proves the values are read. The array is a SCHEDULE and not a KV-head count: its non-zero entries are `1`, the single latent KV head MLA has, while upstream requires `num_attention_heads == num_key_value_heads`, so the array form leaves `num_key_value_heads` at upstream's `None -> num_attention_heads` default. When both spellings are present they are cross-checked on the layer KIND and a clash refuses by name; a per-layer array whose length is not `block_count` refuses by name with the shape found; a non-uniform clamp array refuses, because upstream has ONE `swiglu_limit`; and a file stating neither schedule still refuses, naming both keys. Listed as O18 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) | bug | +| [#2268](https://github.com/mudler/vllm.cpp/issues/2268) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The published GGUF spells MLA geometry by llama.cpp's `attention.key_length` convention and carries no `attention.linear_head_count`, so the loader stops there next.** Measured 2026-08-29 on one tree and one binary, with the [#2243](https://github.com/mudler/vllm.cpp/issues/2243) / [#2177](https://github.com/mudler/vllm.cpp/issues/2177) array fix reverted and restored, driving `LoadedEngine::FromModelDir` at `/mnt/nas_share/rc/ckpt/GLM-5.3-Flash-UD-Q2_K_XL/` on `device = kCPU`, headers only: without the fix it stops at `key glm5next.attention.head_count_kv is not an integer`, with it at `attention.key_length_mla - attention.key_length is -256 but rope.dimension_count is 0` (`glm5_next_weights.cpp:435`). The file is not malformed. llama.cpp writes `key_length = kv_lora_rank + qk_rope_head_dim` and `key_length_mla = qk_nope_head_dim + qk_rope_head_dim` (`b10451:conversion/deepseek.py:345-348`), which for this model gives the artifact's own 512 / 512 / 256 / 256; `scripts/convert-glm5-next-gguf.py` writes `key_length = qk_nope_head_dim`, a different quantity under the same name. `glm5next.attention.linear_head_count`, a `ReqInt` in the builder, is in none of the file's 72 keys and llama.cpp spells it nowhere. Filed rather than fixed in that flow because it moves the WRITE side: putting the reader on llama.cpp's meaning without moving the converter would refuse our own output, and which spelling this project writes is a row-and-spec decision. O7 records that our converter has never been run, so no artifact of ours is invalidated by the move. Recorded under `## Owed` as O18 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) | bug | +| [#2220](https://github.com/mudler/vllm.cpp/issues/2220) | `LTX25-ORACLE-ABSOLUTE` | **A CUDA toolkit staged off CIFS compiles but cannot be LINKED against, and the precondition that should catch it checks the one link that works.** `/workspace` is CIFS and stores no symlink, so a staged toolkit carries only `libcudart.so.13.3.29` and `libcublasLt.so.13.6.0.2`. `ltx25-oracle-absolute-render.sh` rebuilt the links with `b=${f%%.so.*}; ln -sf "$f" "$b.so"; ln -sf "$f" "$b.so.${f#*.so.}"` -- and `${f#*.so.}` strips the SHORTEST prefix, so it expands to `13.3.29` rather than `13`. The second `ln` therefore links `libcudart.so.13.3.29` to ITSELF and **`libcudart.so.13`, the SONAME, is never created**. That is the name `ld` resolves versioned undefined symbols against, so CMake reports `Found CUDAToolkit`, every CUDA TU compiles, and the job dies 21 minutes later with 38 `undefined reference to ...@libcudart.so.13` / `@libcublasLt.so.13` and `ninja: build stopped`. `need_ok` tested `[ -f .../libcublasLt.so ]`, which is exactly the link the loop DID create, so it passed on an unlinkable toolkit -- a precondition that cannot fail. LATENT, not new: the staging branch is a FALLBACK, and every earlier lease found `/usr/local/cuda` 13.0.88 and never took it; `dgx:gpu0` went `unhealthy ... worker_lost` for 3h20m on 2026-08-28 and returned without a toolkit, which exercised it for the first time. The two runs A/B in their own configure logs: `20260827T220845Z` `/usr/local/cuda` 13.0.88 built in 1192 s, `20260828T224529Z` `/root/cudatk` 13.3.73 failed at link. Fixed in flow: take the MAJOR (`v=${f#*.so.}; ${v%%.*}`), prefer `ldconfig -n` which reads each object's own `DT_SONAME`, and assert `.so` resolves AND `.so.` exists for both libraries BEFORE the build. Red-before/green-after on a replica of the CIFS layout: old loop creates no `.so.13`, new logic creates both, and the guard FAILS on the old layout, PASSES on the new, and FAILS on the real NAS source. `rc` job `1ad519b1-4e75-41d7-9386-9932076390f1`, exit 34. Also recorded in [`environment.md`](environment.md) as a lease-environment fact, because it will bite the next row | bug | +| [#2186](https://github.com/mudler/vllm.cpp/issues/2186) | `MODEL-DSV4-EXL3` | **The DeepSeek-V4-Flash EXL3 artifact refused at 108.59 GiB because W1c materialized the carried tower to f32 from one- and two-byte storage, and 26.64 GiB of the total was that widening.** FIXED by W1d: the nine FP8-sourced fields of `DeepseekV4LayerHostWeights` (`wq_a`, `wq_b`, `wkv`, `wo_a`, `wo_b`, `idx_wq`, `shared_w1/w2/w3`) are held at `HostBf16` instead of `std::vector`, taking that half from 21.82 GiB to 10.91 GiB and the artifact's projected residency from 108.59 GiB to ~97.7 GiB against 119.63 GiB physical. **The narrowing is exactly lossless and the gate says so rather than accepting a tolerance**: E4M3 carries four significand bits and E8M0 is a pure power of two, so every value of this tower is exactly representable in bf16's eight, and `narrowing_lost_a_bit == 0` is asserted per element alongside the existing value equality. `Dot`, `MatVec`, `Gemm`, `GroupedOutputLora` and the CUDA `GroupedOLoraKernel` widen each weight AS THEY READ IT -- half the bytes moved, f32 accumulators, reduction order unchanged -- so nothing is materialized back to f32 on any path. **The residency accounting is now gated for the first time**: every other residency case in `test_deepseek_v4_exl3_loader.cpp` compares `DeepseekV4HostResidentBytes` against itself and stays green for any self-consistent formula, so a new case rebuilds the total from each loaded field's own `value_type`; hardcoding `sizeof(float)` back into the accounting takes exactly that one case red (MUTATION-PROVEN 2026-08-29). **Does NOT claim the artifact runs**: the DSA composition is still unported and the forward still refuses by name on the 21 `compress_ratio == 4` layers ([#1961](https://github.com/mudler/vllm.cpp/issues/1961), [#1970](https://github.com/mudler/vllm.cpp/issues/1970), [#1976](https://github.com/mudler/vllm.cpp/issues/1976)), and the ~97.7 GiB figure is a projection from the measured split, not a load that has been observed to complete. The remaining ~2.6 GiB of widening is the BF16-sourced norms/embeddings/router, left at f32 and still owed. Spec [model-dsv4-exl3.md](specs/model-dsv4-exl3.md) `## W1d design` | bug | | [#2244](https://github.com/mudler/vllm.cpp/issues/2244) | `BACKEND-TENSTORRENT-QWEN35` | **The staging upload still pays tt-metal's full creation path on every step.** W4 (#2107) landed levers 1+2 — bulk bf16 staging and single-slot resolution, 0.104 → 0.177 tok/s (+70%), staging `Numel()` share 27.09% → 1.76% — and re-attributed the residual: ~23% of the staging chain is tt-metal per-upload internal work (a fresh `MeshBuffer` allocation, cluster/chip discovery, CQ completion handling) and ~19.2% is CPU threadpool spin, with lever 3 (batch per-layer staging) explicitly not taken. Every upload still goes through `UploadRowsBf16` (`src/vt/tenstorrent/tenstorrent_ops.cpp:469`), which builds a new `ttnn::Tensor` via `from_span`, so identical geometry pays the creation path every step. The lever the W4 record named: allocate the device buffer once per staging slot — lifecycle tied to the slot structures under the #1486 never-destroy rule — and write the host bytes through the mesh command queue (`MeshCommandQueue::enqueue_write`/`enqueue_write_shard`), making the upload allocation-free. The tt-metal-internal half is a proof obligation, not an assumption: read the pinned tt-metal's mesh write path and trace the executed chain before declaring any part of the lever unreachable. `StagingStats` gains route counters for the new path; the capture-unsafe host-write refusals keep their semantics; the f32-conversion arms keep their declared dtypes. Invariant: staging stays bit-identical — the sacred golden pair 16/16 and the full TT suite green; this wave changes speed, never tokens. Evidence owed: same-method before/after profile on the P150 (identical leg, lock discipline) plus a fresh benchmark-record entry; a wall that does not move is a reported result, not a failure — the attribution shifts or the lever is named unreachable with the trace that proves it. Owned by `BACKEND-TENSTORRENT-QWEN35`, next wave after W3 (#2201, landed via #2217) | feature | -| [#2260](https://github.com/mudler/vllm.cpp/issues/2260) | `QUANT-GGUF-IQ2_XS` | **CUDA has no keep-quant kernel for IQ2_XS or IQ4_XS, so the GLM-5.3-Flash artifact FITS `dgx:gpu0` and does not RUN there — the expert GEMM falls back to the CPU, and the fused seam throws.** Found reviewing [#2256](https://github.com/mudler/vllm.cpp/pull/2256), which lands the two CPU keep-quant `vec_dot` kernels and thereby flips the artifact's 82 IQ2_XS and 3 IQ4_XS tensors from `kExpandBf16` to `kKeepQuant`, taking resident cost 426.72 -> 101.14 GiB. `IsCudaKeepQuantSupported` admits ten Q8_K-family encodings (IQ2_XXS, IQ3_XXS, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_S, IQ1_S, IQ1_XXXS) and neither of these two, while `DeviceKeepQuantSupported` returns `true` for CUDA on its `default:` arm regardless, on the recorded ground that CUDA falls back to the CPU kernel for anything it lacks. So `MatmulBTQuantGroupedKernelCuda` round-trips every grouped expert GEMM to the host cores behind a full `cudaStreamSynchronize`, and `MoeGateUpSwiGLUGroupedCuda` THROWS `gate/up must be the SAME CUDA keep-quant dtype` because `MergedGemm` selects the fused op on device registration alone with no dtype predicate. NOT reachable today — `glm5_next_moe.cpp` is W5's host reference and `laguna.cpp` is the only model on the fused seam — so #2256 breaks nothing; it becomes live when AGENTS.md's `vt::MergedGemmGroup` routing lands in W5b ([#2241](https://github.com/mudler/vllm.cpp/issues/2241)) / W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)), and a 101 GiB-resident model then throws at first forward. Three options in the issue: port the two CUDA kernels, keep expanding these two on CUDA (honest, but the artifact stops fitting), or refuse by name at load instead of throwing with the model resident. Owning rows `QUANT-GGUF-IQ2_XS` and `QUANT-GGUF-IQ4_XS` in [`quantization-matrix.md`](../quantization-matrix.md), both carrying the disclosure in place; also carried as **O19** under `## Owed` in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md). Campaign [#1998](https://github.com/mudler/vllm.cpp/issues/1998) | bug | -| [#2277](https://github.com/mudler/vllm.cpp/issues/2277) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The published GLM-5.3-Flash GGUF is `tokenizer.ggml.pre = "glm4"`, and our pre-tokenizer table refuses that name — this is where the loader stops once [#2268](https://github.com/mudler/vllm.cpp/issues/2268) is fixed.** Measured 2026-08-29 on one tree and one build directory, three legs of one probe object driven through `LoadedEngine::FromModelDir` on `device = kCPU` at `/mnt/nas_share/rc/ckpt/GLM-5.3-Flash-UD-Q2_K_XL/`, headers only: the baseline reader stops at `attention.key_length_mla - attention.key_length is -256 but rope.dimension_count is 0`; with the MLA convention fixed but `attention.linear_head_count` still required it stops at `missing metadata key glm5next.attention.linear_head_count`, one key along; with both fixed it stops at `tokenizer: unsupported tokenizer.ggml.pre "glm4"`, past config resolution entirely. `src/vllm/tokenizer/tokenizer.cpp::FromGguf` maps seven pre names — `qwen35`, `qwen2`, `llama-bpe`, the four GPT-4o names, `deepseek-llm`, the three DeepSeek-V3 names and `laguna` — and refuses the rest by name. `glm4` is what every GLM-4 / GLM-5 GGUF carries; shard 1's KV block states `tokenizer.ggml.model = gpt2`, `tokenizer.ggml.pre = glm4`, 154880 tokens and 321649 merges. **The splitting rule is free and the BOS is not.** llama.cpp maps `glm4` and `chatglm-bpe` to `LLAMA_VOCAB_PRE_TYPE_CHATGLM4` (`b10451:src/llama-vocab.cpp:2256-2258`), whose regex at `:398` is BYTE-IDENTICAL to `LLAMA_VOCAB_PRE_TYPE_LLAMA3`'s at `:289`, so `SplitPattern::kLlama3` is EXACT here rather than the "close approximation" that [#347](https://github.com/mudler/vllm.cpp/issues/347) and [#1924](https://github.com/mudler/vllm.cpp/issues/1924) each had to undo — compare the two byte strings in the fix rather than trusting this sentence. But the same branch sets `special_bos_id = LLAMA_TOKEN_NULL` (`:2259`) while the artifact states `tokenizer.ggml.bos_token_id = 154822`, so llama.cpp DISCARDS a BOS id the file carries; a port that reads it and prepends it emits one token no reference run emits, on every request, and a token gate built from our own tokenizer could not see it because both sides would agree. Scope: map both names onto the CHATGLM4 rule with the byte comparison recorded rather than asserted, mirror the `special_bos_id` suppression with a case that fails if a BOS is prepended, and gate through `FromModelDir` on a `pre = "glm4"` fixture so the refusal that moves is the production one. Recorded as O20 in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md), which carries the paired measurement | bug | -| [#2279](https://github.com/mudler/vllm.cpp/issues/2279) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **`FromGguf` never reads `tokenizer.ggml.add_bos_token`, so a GGUF that asks for a leading BOS silently gets none.** Found while adding the `glm4` pre name for [#2277](https://github.com/mudler/vllm.cpp/issues/2277) and deliberately not fixed in that flow: #2277's scope is one pre name, this is a property of every GGUF tokenizer this tree loads. llama.cpp reads the key at `b10451:src/llama-vocab.cpp:2585-2586`, and `add_bos` is the ONLY thing that decides the prepend (`:3382-3384`, `if (add_special && add_bos)`); `tok::Tokenizer::FromGguf` reads `tokenizer.ggml.bos_token_id` and stops there, leaving `template_bos_` at -1 so `EncodeWithSpecialTokens` reduces to `Encode` for every GGUF. Nothing is red today because no artifact this tree gates on states the key -- the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL carries 72 KV entries and it is not among them, parsed 2026-08-29 from shard 1's own KV block, so llama.cpp's `add_bos` stays at its `:1815` default `false` and our silence is the right answer there. It is already live in the other direction on the `llama-bpe` family, whose arm at `:2157-2159` sets `add_bos = true` where the `glm4` arm at `:2256-2259` sets nothing, masked only because that path has never been token-gated against llama.cpp with `add_special = true`. No gate can see this class of defect: a prompt short by exactly one leading token still decodes to fluent text, still has a valid shape, still loads and still generates, and a token gate built from our own tokenizer compares us against us. Scope: read `add_bos_token` (and `add_eos_token`, the same upstream block) defaulting to llama.cpp's `false`; decide what represents it, since `template_bos_` has the right meaning and the wrong provenance comment; a case per arm proving exactly one BOS when true, none when false or absent, both round-tripping; and enumerate which committed fixtures and staged artifacts declare the key so the blast radius is measured rather than assumed. Recorded as O21 in [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) | bug | +| [#2260](https://github.com/mudler/vllm.cpp/issues/2260) | `QUANT-GGUF-IQ2_XS` | **CUDA has no keep-quant kernel for IQ2_XS or IQ4_XS, so the GLM-5.3-Flash artifact FITS `dgx:gpu0` and does not RUN there — the expert GEMM falls back to the CPU, and the fused seam throws.** Found reviewing [#2256](https://github.com/mudler/vllm.cpp/pull/2256), which lands the two CPU keep-quant `vec_dot` kernels and thereby flips the artifact's 82 IQ2_XS and 3 IQ4_XS tensors from `kExpandBf16` to `kKeepQuant`, taking resident cost 426.72 -> 101.14 GiB. `IsCudaKeepQuantSupported` admits ten Q8_K-family encodings (IQ2_XXS, IQ3_XXS, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_S, IQ1_S, IQ1_XXXS) and neither of these two, while `DeviceKeepQuantSupported` returns `true` for CUDA on its `default:` arm regardless, on the recorded ground that CUDA falls back to the CPU kernel for anything it lacks. So `MatmulBTQuantGroupedKernelCuda` round-trips every grouped expert GEMM to the host cores behind a full `cudaStreamSynchronize`, and `MoeGateUpSwiGLUGroupedCuda` THROWS `gate/up must be the SAME CUDA keep-quant dtype` because `MergedGemm` selects the fused op on device registration alone with no dtype predicate. NOT reachable today — `glm5_next_moe.cpp` is W5's host reference and `laguna.cpp` is the only model on the fused seam — so #2256 breaks nothing; it becomes live when AGENTS.md's `vt::MergedGemmGroup` routing lands in W5b ([#2241](https://github.com/mudler/vllm.cpp/issues/2241)) / W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)), and a 101 GiB-resident model then throws at first forward. Three options in the issue: port the two CUDA kernels, keep expanding these two on CUDA (honest, but the artifact stops fitting), or refuse by name at load instead of throwing with the model resident. Owning rows `QUANT-GGUF-IQ2_XS` and `QUANT-GGUF-IQ4_XS` in [`quantization-matrix.md`](quantization-matrix.md), both carrying the disclosure in place; also carried as **O19** under `## Owed` in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md). Campaign [#1998](https://github.com/mudler/vllm.cpp/issues/1998) | bug | +| [#2277](https://github.com/mudler/vllm.cpp/issues/2277) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The published GLM-5.3-Flash GGUF is `tokenizer.ggml.pre = "glm4"`, and our pre-tokenizer table refuses that name — this is where the loader stops once [#2268](https://github.com/mudler/vllm.cpp/issues/2268) is fixed.** Measured 2026-08-29 on one tree and one build directory, three legs of one probe object driven through `LoadedEngine::FromModelDir` on `device = kCPU` at `/mnt/nas_share/rc/ckpt/GLM-5.3-Flash-UD-Q2_K_XL/`, headers only: the baseline reader stops at `attention.key_length_mla - attention.key_length is -256 but rope.dimension_count is 0`; with the MLA convention fixed but `attention.linear_head_count` still required it stops at `missing metadata key glm5next.attention.linear_head_count`, one key along; with both fixed it stops at `tokenizer: unsupported tokenizer.ggml.pre "glm4"`, past config resolution entirely. `src/vllm/tokenizer/tokenizer.cpp::FromGguf` maps seven pre names — `qwen35`, `qwen2`, `llama-bpe`, the four GPT-4o names, `deepseek-llm`, the three DeepSeek-V3 names and `laguna` — and refuses the rest by name. `glm4` is what every GLM-4 / GLM-5 GGUF carries; shard 1's KV block states `tokenizer.ggml.model = gpt2`, `tokenizer.ggml.pre = glm4`, 154880 tokens and 321649 merges. **The splitting rule is free and the BOS is not.** llama.cpp maps `glm4` and `chatglm-bpe` to `LLAMA_VOCAB_PRE_TYPE_CHATGLM4` (`b10451:src/llama-vocab.cpp:2256-2258`), whose regex at `:398` is BYTE-IDENTICAL to `LLAMA_VOCAB_PRE_TYPE_LLAMA3`'s at `:289`, so `SplitPattern::kLlama3` is EXACT here rather than the "close approximation" that [#347](https://github.com/mudler/vllm.cpp/issues/347) and [#1924](https://github.com/mudler/vllm.cpp/issues/1924) each had to undo — compare the two byte strings in the fix rather than trusting this sentence. But the same branch sets `special_bos_id = LLAMA_TOKEN_NULL` (`:2259`) while the artifact states `tokenizer.ggml.bos_token_id = 154822`, so llama.cpp DISCARDS a BOS id the file carries; a port that reads it and prepends it emits one token no reference run emits, on every request, and a token gate built from our own tokenizer could not see it because both sides would agree. Scope: map both names onto the CHATGLM4 rule with the byte comparison recorded rather than asserted, mirror the `special_bos_id` suppression with a case that fails if a BOS is prepended, and gate through `FromModelDir` on a `pre = "glm4"` fixture so the refusal that moves is the production one. Recorded as O20 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md), which carries the paired measurement | bug | +| [#2279](https://github.com/mudler/vllm.cpp/issues/2279) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **`FromGguf` never reads `tokenizer.ggml.add_bos_token`, so a GGUF that asks for a leading BOS silently gets none.** Found while adding the `glm4` pre name for [#2277](https://github.com/mudler/vllm.cpp/issues/2277) and deliberately not fixed in that flow: #2277's scope is one pre name, this is a property of every GGUF tokenizer this tree loads. llama.cpp reads the key at `b10451:src/llama-vocab.cpp:2585-2586`, and `add_bos` is the ONLY thing that decides the prepend (`:3382-3384`, `if (add_special && add_bos)`); `tok::Tokenizer::FromGguf` reads `tokenizer.ggml.bos_token_id` and stops there, leaving `template_bos_` at -1 so `EncodeWithSpecialTokens` reduces to `Encode` for every GGUF. Nothing is red today because no artifact this tree gates on states the key -- the staged `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL carries 72 KV entries and it is not among them, parsed 2026-08-29 from shard 1's own KV block, so llama.cpp's `add_bos` stays at its `:1815` default `false` and our silence is the right answer there. It is already live in the other direction on the `llama-bpe` family, whose arm at `:2157-2159` sets `add_bos = true` where the `glm4` arm at `:2256-2259` sets nothing, masked only because that path has never been token-gated against llama.cpp with `add_special = true`. No gate can see this class of defect: a prompt short by exactly one leading token still decodes to fluent text, still has a valid shape, still loads and still generates, and a token gate built from our own tokenizer compares us against us. Scope: read `add_bos_token` (and `add_eos_token`, the same upstream block) defaulting to llama.cpp's `false`; decide what represents it, since `template_bos_` has the right meaning and the wrong provenance comment; a case per arm proving exactly one BOS when true, none when false or absent, both round-tripping; and enumerate which committed fixtures and staged artifacts declare the key so the blast radius is measured rather than assumed. Recorded as O21 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) | bug | | [#2273](https://github.com/mudler/vllm.cpp/issues/2273) | `BACKEND-TENSTORRENT-QWEN35` | **The wall is the per-CQ-operation tt-metal stack, charged once per staging write, and a decode step pays it once per staged tensor.** W5 (#2244) made uploads allocation-free and the wall honestly did not move (−0.14%, noise), and the trace split W4's hypothesis: `allocate_mesh_tensor_on_device_with_topology` is 0.02% of the AFTER profile and the write stacks are identical in both arms, so allocation was never the wall. What remains is the fixed per-op tax — `Threadpool::PollForWork` 14.29%, `MetalContext::instance` 11.14%, `memcpy` 6.23%, `Cluster::get_chip` 5.90%, `read_cq_host_ptr` 5.27% + sub-slices — multiplied by the layer fan-in. The lever (W4 lever 3, deferred there as optional, re-derived as owed): batch per-layer staging — pack a step's staged host rows into one contiguous host block and issue ONE mesh-CQ write per step or layer group, so the per-op tax divides by the fan-in. Invariant: staging stays bit-identical — the sacred golden pair 16/16 STRICT and the full TT suite green; this wave changes speed, never tokens. `StagingStats` gains route counters for the new path; capture-unsafe host-write refusals keep their semantics; f32-conversion arms keep their declared dtypes; a batched/arena layout must state its restage semantics explicitly (W5 review aliasing awareness — same-geometry restage aliases the persistent buffer), and the route must be production-reachable, not test-only. Evidence owed: same-method before/after profile on the P150 (identical leg, JIT-discard per arm, one lock hold) plus a fresh benchmark-record entry; the attribution shifts or the lever is named unreachable with the trace that proves it. The tt-metal-side residual (cached context handles, amortized CQ polling) stays recorded as the upstream-shaped alternative. Owned by `BACKEND-TENSTORRENT-QWEN35`, next wave after W5 (#2244, landed via #2258) | perf | -| [#2286](https://github.com/mudler/vllm.cpp/issues/2286) | `MODEL-DSV4-DSA-COMPOSE` | **The DeepSeek-V4 DSA composition had no owning row, and the forward's own refusal message said so** (`src/vllm/model_executor/models/deepseek_v4.cpp:~738`: "The DSA port itself is OWED and has no owning row"). SCOPED 2026-08-29 by [dsv4-dsa-compose.md](../specs/dsv4-dsa-compose.md), read at the parity pin `5559679229`. It is the blocker between a DeepSeek-V4-Flash artifact that LOADS ([#2186](https://github.com/mudler/vllm.cpp/issues/2186)/[#2283](https://github.com/mudler/vllm.cpp/issues/2283)) and one that RUNS. **The two kernel primitives already have rows** (`KERNEL-ATTN-DSA-SPARSE-INDEX`, `KERNEL-ATTN-DSA-COMPRESSOR`, both `SPIKE`); what had no owner is the ASSEMBLY into `AttentionBlock` -- three layer shapes selected by `compress_ratio` (`attention.py:454-533`), all ending at `forward_mqa` then `_o_proj`. **THREE FINDINGS THE SCOPING ADDS.** (1) The 3-way stream overlap is PERFORMANCE, not correctness: `attention_impl` dispatches with `enable=aux_streams is not None` and ROCm runs the same work sequentially, so a sequential first wave MIRRORS upstream -- stated so a later reader does not "restore" the overlap believing correctness depended on it. (2) `coff = 1 + (compress_ratio == 4)` is a per-token ROLE selected by offset within the gathering window -- `head_offset = (tokens >= COMPRESS_RATIO) * HEAD_SIZE`, emission at boundary tokens only, the state cache holding TWO head-sized rows per token, so a token in the overlap has a DIFFERENT role in each of the two windows containing it. That is the whole of what "never recoverable from the tensor alone" means, and our loader already materializes the width correctly ([#1970](https://github.com/mudler/vllm.cpp/issues/1970)), making this a FORWARD change rather than a loader one. (3) The compressor is two stages with the second boundary-gated, and its RoPE is exactly specified -- GPT-J style, `is_neox_style=False` (interleaved, NOT split-half), on the LAST `rope_head_dim` elements, at position `(positions // compress_ratio) * compress_ratio`. **HARD ORDERING:** consumes `KV-DSV4-MULTICACHE` ([#1925](https://github.com/mudler/vllm.cpp/issues/1925)) and W1 cannot start before that row's W3 hands the forward the cache. **NOT GATEABLE AT OR BELOW 512 TOKENS**, because the one arm that caches today forces indexer and compressor OFF and is exact only while `seq_len <= index_topk` (=512) -- so every gate on this row must exceed it. Also records that `config.json`'s `compress_ratios` has **46** entries `{0:5, 4:21, 128:20}` while older records say "43 layers"; 43 is the trellis shard count, and W1 reconciles which number each claim means before writing code. NOT FIXED IN FLOW and deliberately: `AGENTS.md` requires the spec first, and a capability of this size needs agreement on scope before implementation waves start | enhancement | -| [#2283](https://github.com/mudler/vllm.cpp/issues/2283) | `MODEL-DSV4-EXL3` | **The DeepSeek-V4 carried tower's BF16-sourced half is still widened to f32 (~2.62 GiB), and W1d's ~97.7 GiB projection has never been observed.** Filed 2026-08-29 because W1d ([#2186](https://github.com/mudler/vllm.cpp/issues/2186), landed `c9ad53fee`) CLOSED its issue while `.agents/specs/model-dsv4-exl3.md` `## Owed` still pointed two live entries at it -- a reader following either landed on a closed issue. **(1)** The 108.59 -> ~97.7 GiB figure is arithmetic on the measured 108.59 / 26.64 split, not a load anyone has watched complete; the last real measurement (2026-08-28, `dgx:gpu0`, worker `rc-worker-4b8lj`, tree `525d2b991`) REFUSED, and nothing has re-run since. It falls due as an `rc`-leased `dgx:gpu0` measurement against the staged 100 GB artifact, and a load that completes is still not a forward that runs (#1961, #1970, #1976). **(2)** The carried tower's other half -- norms, embeddings, router, `BF16` on disk, 2.621 GiB -> 5.24 GiB at f32 -- is untouched, and the same "Inherit vLLM defaults" argument applies verbatim. NOT folded into W1d deliberately: W1d's nine fields had three consuming functions and one device vtable entry, while this half is read by the sampler and lm_head paths too (the embedding is held twice on device, #1946), so it is a wave with its own gate. Build on what W1d left: `vllm::HostBf16`, the inlined `vllm::HostBf16ToF32` (out-of-line `vt::BF16ToF32` + no LTO would cost a call per element in the GEMV inner loop) with its exhaustive 65536-pattern agreement case, `Dot`'s bf16 overload, generic `MatVec`/`Gemm`/`expert_f32`/`GroupedOutputLora`, and a `DeepseekV4HostResidentBytes` that now reads each field's own `value_type` under a mutation-proven gate. Spec [model-dsv4-exl3.md](../specs/model-dsv4-exl3.md) `## Owed` | bug | +| [#2283](https://github.com/mudler/vllm.cpp/issues/2283) | `MODEL-DSV4-EXL3` | **The DeepSeek-V4 carried tower's BF16-sourced half is still widened to f32 (~2.62 GiB), and W1d's ~97.7 GiB projection has never been observed.** Filed 2026-08-29 because W1d ([#2186](https://github.com/mudler/vllm.cpp/issues/2186), landed `c9ad53fee`) CLOSED its issue while `.agents/specs/model-dsv4-exl3.md` `## Owed` still pointed two live entries at it -- a reader following either landed on a closed issue. **(1)** The 108.59 -> ~97.7 GiB figure is arithmetic on the measured 108.59 / 26.64 split, not a load anyone has watched complete; the last real measurement (2026-08-28, `dgx:gpu0`, worker `rc-worker-4b8lj`, tree `525d2b991`) REFUSED, and nothing has re-run since. It falls due as an `rc`-leased `dgx:gpu0` measurement against the staged 100 GB artifact, and a load that completes is still not a forward that runs (#1961, #1970, #1976). **(2)** The carried tower's other half -- norms, embeddings, router, `BF16` on disk, 2.621 GiB -> 5.24 GiB at f32 -- is untouched, and the same "Inherit vLLM defaults" argument applies verbatim. NOT folded into W1d deliberately: W1d's nine fields had three consuming functions and one device vtable entry, while this half is read by the sampler and lm_head paths too (the embedding is held twice on device, #1946), so it is a wave with its own gate. Build on what W1d left: `vllm::HostBf16`, the inlined `vllm::HostBf16ToF32` (out-of-line `vt::BF16ToF32` + no LTO would cost a call per element in the GEMV inner loop) with its exhaustive 65536-pattern agreement case, `Dot`'s bf16 overload, generic `MatVec`/`Gemm`/`expert_f32`/`GroupedOutputLora`, and a `DeepseekV4HostResidentBytes` that now reads each field's own `value_type` under a mutation-proven gate. Spec [model-dsv4-exl3.md](specs/model-dsv4-exl3.md) `## Owed` | bug | +| [#2286](https://github.com/mudler/vllm.cpp/issues/2286) | `MODEL-DSV4-DSA-COMPOSE` | **The DeepSeek-V4 DSA composition had no owning row, and the forward's own refusal message said so** (`src/vllm/model_executor/models/deepseek_v4.cpp:~738`: "The DSA port itself is OWED and has no owning row"). SCOPED 2026-08-29 by [dsv4-dsa-compose.md](specs/dsv4-dsa-compose.md), read at the parity pin `5559679229`. It is the blocker between a DeepSeek-V4-Flash artifact that LOADS ([#2186](https://github.com/mudler/vllm.cpp/issues/2186)/[#2283](https://github.com/mudler/vllm.cpp/issues/2283)) and one that RUNS. **The two kernel primitives already have rows** (`KERNEL-ATTN-DSA-SPARSE-INDEX`, `KERNEL-ATTN-DSA-COMPRESSOR`, both `SPIKE`); what had no owner is the ASSEMBLY into `AttentionBlock` -- three layer shapes selected by `compress_ratio` (`attention.py:454-533`), all ending at `forward_mqa` then `_o_proj`. **THREE FINDINGS THE SCOPING ADDS.** (1) The 3-way stream overlap is PERFORMANCE, not correctness: `attention_impl` dispatches with `enable=aux_streams is not None` and ROCm runs the same work sequentially, so a sequential first wave MIRRORS upstream -- stated so a later reader does not "restore" the overlap believing correctness depended on it. (2) `coff = 1 + (compress_ratio == 4)` is a per-token ROLE selected by offset within the gathering window -- `head_offset = (tokens >= COMPRESS_RATIO) * HEAD_SIZE`, emission at boundary tokens only, the state cache holding TWO head-sized rows per token, so a token in the overlap has a DIFFERENT role in each of the two windows containing it. That is the whole of what "never recoverable from the tensor alone" means, and our loader already materializes the width correctly ([#1970](https://github.com/mudler/vllm.cpp/issues/1970)), making this a FORWARD change rather than a loader one. (3) The compressor is two stages with the second boundary-gated, and its RoPE is exactly specified -- GPT-J style, `is_neox_style=False` (interleaved, NOT split-half), on the LAST `rope_head_dim` elements, at position `(positions // compress_ratio) * compress_ratio`. **HARD ORDERING:** consumes `KV-DSV4-MULTICACHE` ([#1925](https://github.com/mudler/vllm.cpp/issues/1925)) and W1 cannot start before that row's W3 hands the forward the cache. **NOT GATEABLE AT OR BELOW 512 TOKENS**, because the one arm that caches today forces indexer and compressor OFF and is exact only while `seq_len <= index_topk` (=512) -- so every gate on this row must exceed it. Also records that `config.json`'s `compress_ratios` has **46** entries `{0:5, 4:21, 128:20}` while older records say "43 layers"; 43 is the trellis shard count, and W1 reconciles which number each claim means before writing code. NOT FIXED IN FLOW and deliberately: `AGENTS.md` requires the spec first, and a capability of this size needs agreement on scope before implementation waves start | enhancement | | [#2249](https://github.com/mudler/vllm.cpp/issues/2249) | `MODEL-MM-QWEN4-EXP` | **The interleaved-mRoPE cos\|sin table builder was `static`, so the `qwen4_exp` QSA half of the layer loop could only have a SECOND copy of it.** Item 5 of five prerequisites measured while attempting the loop ([#2031](https://github.com/mudler/vllm.cpp/issues/2031)); W5d-2 closes that item only, and the other four stay open on this issue. `BuildMropeCosSinHost` sat at `src/vllm/model_executor/models/qwen3_5.cpp:9472` with internal linkage and all three of its uses inside that translation unit, and mRoPE is the arithmetic where a duplicate diverges in silence — a wrong axis still produces plausible tokens. **Fixed by `include/vllm/model_executor/models/qwen3_5_mrope.h`**, which takes the SIMPLER of the two shapes this row has already used: `RunGdnBlockPaged` (#2110) and `RunMoeBlock` needed a public WRAPPER because their signatures name `StepDevInputs`, a type qwen3_5.cpp declares privately, while this signature names only `std::vector`, `int64_t` and `vllm::HfConfig`, so the whole extraction is the `static` keyword plus a declaration. The definition does not move: `sed -n '9473,9514p'` of the base-SHA file and of the head both sha256 to `259b1b932cae0611...`. **A byte-identical body is not on its own a value guarantee**, because the keyword that changed is exactly the one deciding which definition a caller binds to, so `tests/vllm/models/test_qwen3_5_mrope.cpp` pins 152 f32 BIT PATTERNS over four cases — interleaved and chunked at the same config and positions, T == 1 at another rotary_dim and rope_theta, and a small-`t`-section case where the `pair <= 3 * sec[k]` boundary decides differently — against what the FILE-STATIC produced at base `94de63ff5`, captured by compiling its `sed`-extracted text in a standalone harness. Bitwise and not an epsilon: a pure host computation over `std::cos`/`std::pow` has no reduction-order freedom, so a tolerance would hide the only defect an extraction can introduce. 26 pre-existing qwen3_5/qwen4_exp suites are identical in exit status and in case and assertion counts before and after (base `94de63ff5` against branch head `c1ccbac19`, both of which predate this branch's merge of `main`; that merge brings W5b's `test_qwen4_exp_forward`, which makes the same glob match 27 on the merged head and is NOT part of this pair, having existed at neither end of it) — but **FOUR of the 26 measure nothing on a host without the checkpoints**, and only `test_qwen35_paged_engine` (rc 77) says so: `test_qwen35_gguf_spec_decode` (3 cases / **0 assertions**, `SKIP: set VLLM_MTP_GGUF_MODEL`), `test_qwen3_5_vl_e2e` and `test_qwen3_5_vl_video_e2e` (1 case / **0 assertions** each, `SKIP: Qwen3.6-27B checkpoint absent`) each exit 0 and print `Status: SUCCESS!`. Those last two are the STRICT token-exact e2e gates on `VLGenerateCoreGdn`, the driver core holding the call sites the reachability mutation deletes, so on such a host the reachability evidence rests ENTIRELY on `test_qwen3_5_moe_vision` (7 cases / 38 assertions, of which one case reds). **One equivalent mutant is recorded rather than hidden:** upstream's `<=` in `mrope.py:60-63` cannot be told from a `<` here, because the guard already requires `pair % 3 == 1` while `3 * sec[k]` is divisible by 3 — the boundary is unreachable, and the mutation that DOES red it shifts the bound instead. Reachability proven by deleting both production call sites, which reds `test_qwen3_5_moe_vision`'s `..._uses_MRoPE_positions_not_plain_1d`; the new suite stays green under that deletion and says so in its own comment, because a unit case measures the function and never that anything reaches it | feature | | [#2257](https://github.com/mudler/vllm.cpp/issues/2257) | `ENG-MM-QWEN36-VL-FORWARD` | **The four Qwen3.5/3.6 VL greedy drivers have no production caller: `ModelRegistry::Forward` cannot route an image or video request to any of them.** `Qwen3_5VLGenerateGreedy`, `Qwen3_5VLGenerateGreedyVideo`, `Qwen3_5MoeVLGenerateGreedy` and `Qwen3_5MoeVLGenerateGreedyVideo` are DEFINED at `src/vllm/model_executor/models/qwen3_5.cpp:9892,9915,9960,9974` and declared in `qwen3_5.h` / `qwen3_5_dense.h`; a grep for the four names over `src/ include/ examples/ tools/ benchmarks/` returns those four definitions and their six declaration lines and NOTHING else, so **every caller is in `tests/`**. The registered factories for `Qwen3_5ForConditionalGeneration` and `Qwen3_5MoeForConditionalGeneration` (`REGISTER_VLLM_MODEL`, `qwen3_5_dense.cpp:283`) route the forward to `ForwardQwen3_5Dense`, which takes a `ModelForwardInput` and carries no multimodal hook, and `ModelRegistry::Forward` additionally refuses a non-null `multi_kv` (`model_registry.cpp:428-440`) that this architecture's three cache groups make the runner set. The tree already states the same condition for the sibling 4B driver at `include/vllm/entrypoints/openai/chat_mm.h:266-267` — the M2c driver "runs it standalone, outside `ModelRegistry::Forward`". So M3-b image and M3d video are gated e2e and correct, and **no user arrives at them**, which by AGENTS.md `## Nothing lands dead` makes every change inside `VLGenerateCoreGdn` or below it reached by a test and by nothing else. FOUND, not caused, while landing W5d-2 of [#2249](https://github.com/mudler/vllm.cpp/issues/2249), which gave `BuildMropeCosSinHost` external linkage: that wave's `## Owed` entry in `.agents/specs/qwen4-exp-flash-next.md` has to name who owns the hop above its call sites, and nothing tracked this gap. The condition PREDATES the extraction and is unchanged by it in either direction. Owned by `ENG-MM-QWEN36-VL-FORWARD`, which owns `BuildMropeCosSinHost`, the shared `VLGenerateCoreGdn` and the two 27B dense drivers; the two MoE drivers additionally sit under `MODEL-MM-qwen3-5-qwen3-5-moe-for-conditional-generation` and [#891](https://github.com/mudler/vllm.cpp/issues/891) | bug | -| [#2242](https://github.com/mudler/vllm.cpp/issues/2242) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **W5c — the weight tower and `load_weights`: `Glm5NextForConditionalGeneration` LOADS.** The GGUF arm of the registry's `load_weights` hook now returns a real `Glm5NextLoadedModel` built by `LoadGlm5NextFromGguf`, so this architecture has a `LoadedModel` for the first time and the loader's refusal is gone from product output. The tower covers every tensor group the architecture declares — the KDA layer with its three separate depthwise convs, the NoPE MLA with the two SPLIT absorbed halves, the DSA k-pool indexer, the flat mHC pair at `(2 + hc_mult) * hc_mult`, the 288 stacked routed experts plus one shared, and the dense MLP on the leading three layers — and refuses BY NAME on a missing tensor, a shape disagreement or a non-negative `ssm_a`. Gated against the REAL published artifact with no asset: `tests/vllm/models/glm5_next_gguf_manifest.inc` freezes the 1412-tensor header table of `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL at revision `d425e572fb9686125831f476129e51cea34bc5b4`, and the name map is accounted against it in BOTH directions — 1383 enumerated, 0 missing, 0 unexplained, 29 MTP-block tensors deliberately dropped, 1383 + 29 = 1412. `blk.45` is NOT built as a decoder layer, asserted three ways because each alone is satisfiable by a wrong loader: no `blk.45.*` name is enumerated, the file demonstrably HAS one, and the loader positively COUNTS the 29 tensors it skipped. Driven at the staged artifact through the same chain `LoadedEngine::FromModelDir` uses, headers only: all four shards open, the config resolves to 45 layers / 34 KDA / 11 DSA / hc_mult 4 / kpool 4 / NoPE, and every one of the 1383 names resolves at 41 MB peak RSS with no payload byte read. The residency the load would take, predicted by `PeekRoute` over those same names: 736 tensors keep their blocks at 98.260 GiB, 647 expand to bf16 at 0.446 GiB. NOT DONE HERE and named in the spec: the forward (W5b, [#2241](https://github.com/mudler/vllm.cpp/issues/2241)), the KV-cache spec, the vision tower (W6), the MTP head (O2) and the safetensors arm, all five still refusing by name. Campaign [#1998](https://github.com/mudler/vllm.cpp/issues/1998), spec [`specs/glm5-next-flash.md`](../specs/glm5-next-flash.md) §W5c | feature | +| [#2242](https://github.com/mudler/vllm.cpp/issues/2242) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **W5c — the weight tower and `load_weights`: `Glm5NextForConditionalGeneration` LOADS.** The GGUF arm of the registry's `load_weights` hook now returns a real `Glm5NextLoadedModel` built by `LoadGlm5NextFromGguf`, so this architecture has a `LoadedModel` for the first time and the loader's refusal is gone from product output. The tower covers every tensor group the architecture declares — the KDA layer with its three separate depthwise convs, the NoPE MLA with the two SPLIT absorbed halves, the DSA k-pool indexer, the flat mHC pair at `(2 + hc_mult) * hc_mult`, the 288 stacked routed experts plus one shared, and the dense MLP on the leading three layers — and refuses BY NAME on a missing tensor, a shape disagreement or a non-negative `ssm_a`. Gated against the REAL published artifact with no asset: `tests/vllm/models/glm5_next_gguf_manifest.inc` freezes the 1412-tensor header table of `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL at revision `d425e572fb9686125831f476129e51cea34bc5b4`, and the name map is accounted against it in BOTH directions — 1383 enumerated, 0 missing, 0 unexplained, 29 MTP-block tensors deliberately dropped, 1383 + 29 = 1412. `blk.45` is NOT built as a decoder layer, asserted three ways because each alone is satisfiable by a wrong loader: no `blk.45.*` name is enumerated, the file demonstrably HAS one, and the loader positively COUNTS the 29 tensors it skipped. Driven at the staged artifact through the same chain `LoadedEngine::FromModelDir` uses, headers only: all four shards open, the config resolves to 45 layers / 34 KDA / 11 DSA / hc_mult 4 / kpool 4 / NoPE, and every one of the 1383 names resolves at 41 MB peak RSS with no payload byte read. The residency the load would take, predicted by `PeekRoute` over those same names: 736 tensors keep their blocks at 98.260 GiB, 647 expand to bf16 at 0.446 GiB. NOT DONE HERE and named in the spec: the forward (W5b, [#2241](https://github.com/mudler/vllm.cpp/issues/2241)), the KV-cache spec, the vision tower (W6), the MTP head (O2) and the safetensors arm, all five still refusing by name. Campaign [#1998](https://github.com/mudler/vllm.cpp/issues/1998), spec [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md) §W5c | feature | | [#2291](https://github.com/mudler/vllm.cpp/issues/2291) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The W7a converter and the published artifact disagree on three tensors, and one of the three is a silent value transform.** Found while implementing W5c ([#2242](https://github.com/mudler/vllm.cpp/issues/2242)), whose own scope sentence assumed they agreed, and fixed in the same flow. Read at source from llama.cpp PR [#27752](https://github.com/ggml-org/llama.cpp/pull/27752) head `8a8d0bcc4d5fdf024c457526245bec4bc3a12adc` (`conversion/glm5next.py`, sha256 `bfacba27746096e7bb3ca4a2549c9026d3475e226c7f3edf230c37ffadc7b6b3`) plus the `DeepseekV2Model` it inherits, and confirmed against the staged UD-Q2_K_XL header table. (1) `.dt_bias` is RENAMED to `.dt_proj.bias` before the generic map runs, so the file carries `blk.N.ssm_dt.bias` and no bare `ssm_dt`. (2) `kv_b_proj` is SPLIT into `attn_k_b` and `attn_v_b` with the k half TRANSPOSED, so the file carries two tensors at DIFFERENT shapes — ne `[256, 512, 64]` and `[512, 256, 64]` — and no `attn_kv_b.weight`; because `qk_nope_head_dim == v_head_dim == 256`, a fixture at equal head dims cannot tell a correct split from a swapped one, so the gate asserts both SHAPES and the nearer-own-half property rather than sizes. (3) `ssm_a` holds `-exp(A_log)`, not `A_log` — the dangerous one, because the tensor is present, the shape is right and the values are plausible floats, so nothing structural fires: a loader that inverts gets NaN on every KDA decay, one that does not runs a sign-flipped forget gate and generates fluent wrong text, and no oracle for this model runs on any device this project reaches to tell the difference. Fixed in the converter, in the C++ name map (the split needs its own 1:1 table, since one HF name maps to two GGUF names and a dict cannot carry one key twice) and in the new loader, which refuses a non-negative `ssm_a` by name. `tests/scripts/test_convert_glm5_next_gguf.py` was RED on the tensor set before the converter moved | bug | -| [#2300](https://github.com/mudler/vllm.cpp/issues/2300) | `ENG-MM-INPUT-PIPELINE` | **The GPU runner never sets `ModelForwardInput.mm`, so a Qwen3-VL server throws on the first forward step of every request, text or image.** Measured at `e541be98a`. `ForwardQwen3VLForConditionalGeneration` opens with `VT_CHECK(input.mm.has_value(), ...)` at `src/vllm/model_executor/models/qwen3_vl_registry.cpp:127`, and that forward is what a loaded Qwen3-VL resolves to: `REGISTER_VLLM_MODEL(qwen3_vl, "Qwen3VLForConditionalGeneration", ...)` at `:203` binds `.forward` at `:184`. The field it demands is `std::optional mm = std::nullopt` (`include/vllm/model_executor/models/model_registry.h:446`), and the runner's designated initializer at `src/vllm/v1/worker/gpu/runner.cpp:2234` names 16 fields and NOT `.mm` before calling `ModelRegistry::Forward` at `:2340`. Over the whole 4443-line file a grep for `mm_features`, `MultiModalForwardInput` or `.mm = ` returns 0, and so does a grep for `mm`, `multimodal` or `MultiModal`; `include/vllm/v1/worker/gpu/input_batch.h:90` records the worker input batch as a subset with "mm_features / generator / lora / prompt_embeds / pooling DEFERRED", so the features the field would be built from never reach the worker at all. The three writers of `.mm` in the tree are single-sequence drivers (`qwen3_vl.cpp:638`, `gemma4_mm.cpp:240`, `muse_glimmer_mm.cpp:358`), none of them the runner. **The refusing shape is a per-model choice, not a tree-wide one:** `gemma4_registry.cpp:145` and `muse_glimmer_registry.cpp:113` both guard with `if (input.mm.has_value())` and both carry the sentence "nullopt on every text step => the text path below is byte-identical", so under those two a runner step with `mm` unset runs the text path while Qwen3-VL throws. **STATICALLY DERIVED and NOT RUN:** no binary was built and no server was started, because the filing unit touches no product code; every claim is a `file:line` read plus the two grep counts, re-derived at the base SHA, and a runtime confirmation needs real safetensors weights because `LoadQwen3VLForConditionalGeneration` (`qwen3_vl_registry.cpp:95`) refuses any other source. Distinct from [#1358](https://github.com/mudler/vllm.cpp/issues/1358), which is the same root cause with a different symptom (the tower is loaded on the production path and read by nothing, costing memory rather than every request), and from [#2257](https://github.com/mudler/vllm.cpp/issues/2257), which is the Qwen3.5 and Qwen3.6 VL drivers having no production caller at all (`ForwardQwen3_5Dense` carries no multimodal hook, so it never reads `input.mm` and never throws). NOT fixed in flow: the repair is either a text arm in the registered forward or the runner building `mm` from staged encoder outputs, and the choice between them changes what an image request does, so it takes the surprising-fix path with its own spec, gate and independent review rather than an in-flow repair. Owned by `ENG-MM-INPUT-PIPELINE`, listed under `## Owed` in [`multimodal-track.md`](../specs/multimodal-track.md), and corrected on [`mm-serving.md`](../specs/mm-serving.md) | bug | +| [#2300](https://github.com/mudler/vllm.cpp/issues/2300) | `ENG-MM-INPUT-PIPELINE` | **The GPU runner never sets `ModelForwardInput.mm`, so a Qwen3-VL server throws on the first forward step of every request, text or image.** Measured at `e541be98a`. `ForwardQwen3VLForConditionalGeneration` opens with `VT_CHECK(input.mm.has_value(), ...)` at `src/vllm/model_executor/models/qwen3_vl_registry.cpp:127`, and that forward is what a loaded Qwen3-VL resolves to: `REGISTER_VLLM_MODEL(qwen3_vl, "Qwen3VLForConditionalGeneration", ...)` at `:203` binds `.forward` at `:184`. The field it demands is `std::optional mm = std::nullopt` (`include/vllm/model_executor/models/model_registry.h:446`), and the runner's designated initializer at `src/vllm/v1/worker/gpu/runner.cpp:2234` names 16 fields and NOT `.mm` before calling `ModelRegistry::Forward` at `:2340`. Over the whole 4443-line file a grep for `mm_features`, `MultiModalForwardInput` or `.mm = ` returns 0, and so does a grep for `mm`, `multimodal` or `MultiModal`; `include/vllm/v1/worker/gpu/input_batch.h:90` records the worker input batch as a subset with "mm_features / generator / lora / prompt_embeds / pooling DEFERRED", so the features the field would be built from never reach the worker at all. The three writers of `.mm` in the tree are single-sequence drivers (`qwen3_vl.cpp:638`, `gemma4_mm.cpp:240`, `muse_glimmer_mm.cpp:358`), none of them the runner. **The refusing shape is a per-model choice, not a tree-wide one:** `gemma4_registry.cpp:145` and `muse_glimmer_registry.cpp:113` both guard with `if (input.mm.has_value())` and both carry the sentence "nullopt on every text step => the text path below is byte-identical", so under those two a runner step with `mm` unset runs the text path while Qwen3-VL throws. **STATICALLY DERIVED and NOT RUN:** no binary was built and no server was started, because the filing unit touches no product code; every claim is a `file:line` read plus the two grep counts, re-derived at the base SHA, and a runtime confirmation needs real safetensors weights because `LoadQwen3VLForConditionalGeneration` (`qwen3_vl_registry.cpp:95`) refuses any other source. Distinct from [#1358](https://github.com/mudler/vllm.cpp/issues/1358), which is the same root cause with a different symptom (the tower is loaded on the production path and read by nothing, costing memory rather than every request), and from [#2257](https://github.com/mudler/vllm.cpp/issues/2257), which is the Qwen3.5 and Qwen3.6 VL drivers having no production caller at all (`ForwardQwen3_5Dense` carries no multimodal hook, so it never reads `input.mm` and never throws). NOT fixed in flow: the repair is either a text arm in the registered forward or the runner building `mm` from staged encoder outputs, and the choice between them changes what an image request does, so it takes the surprising-fix path with its own spec, gate and independent review rather than an in-flow repair. Owned by `ENG-MM-INPUT-PIPELINE`, listed under `## Owed` in [`multimodal-track.md`](specs/multimodal-track.md), and corrected on [`mm-serving.md`](specs/mm-serving.md) | bug | | [#2274](https://github.com/mudler/vllm.cpp/issues/2274) | `SPEC-DFLASH2` | **The DFlash2 paged draft route reads out of bounds EAGERLY at `max_num_seqs=1`, so the committed speed gate cannot measure our arm at all.** `scripts/dflash2-speed-gate.sh` on `main` refuses with `GATE_RC=2 / RESULT_PRESENT=no`: `vllm-cli` exits 134 with `vt cuda: cudaMemcpyAsync: an illegal memory access`, while the oracle arm completes in the same run. Reproducer needs no concurrency, no CUDA graph and no FA2 lane: `VT_DFLASH_PAGED=1 VT_DFLASH_GRAPH=0 vllm-cli --prompt "The capital of France is" --max-tokens 64 --repeat 5 --max-num-seqs 1`. **Repetition 1 completes 64 tokens and a LATER one dies**, so it needs state carried across requests — which is why 16-token single-shot probes survive. SEVEN candidates tested and excluded, each on an FA2-carrying build on one boot: the CUDA graph (`VT_DFLASH_GRAPH=0` still faults, so every earlier `cudaGraphLaunch` attribution was incidental), the FA2 block lane (`VT_FA2_DFLASH_BLOCK=0`), merged QKV (`VT_QWEN3_QKV_MERGE=0` plus an ON control on the same boot), **the whole seam adoption of [#2207](https://github.com/mudler/vllm.cpp/issues/2207) by building `c9b2049bc~1` = `f01fcc4c6`, which still faults and so exonerates it**, FA2 being compiled out (four earlier gate runs were measured on a binary with an EMPTY `CUDA FA2 compiled-arch manifest` because the staged lease script omitted `-DVLLM_CPP_CUTLASS_FETCH=ON` — the same defect that forced the 2026-08-24 retraction), the `max_seq_len` replay staleness fixed in `41dd3398a`, and the host-side bounds accounting (`ddd527f3f` added four shape-only checks that run on EVERY backend; they are silent on the failing configuration, so that class is excluded and the checks remain as a named refusal for whoever breaks the accounting later). `VT_DFLASH_PAGED=0` is the ONLY configuration that completes, and on it the gate PASSES: `GATE_RC=0`, **ours 12.361 tok/s vs vLLM 16.292, ratio 0.759 — 24% slower** at 0.789% SM-clock spread on one boot, which makes this issue NECESSARY BUT NOT SUFFICIENT for parity. `compute-sanitizer` cannot see the fault: the `vt cuda drop-in` layer initialises CUDA before the sanitizer interposes, so memcheck disables itself and reports THAT as its own "1 error" — two leases spent learning it, recorded so a third is not. The remaining lead is the one class the pre-existing guards cannot check on CUDA, because both are `kCPU`-guarded and their own comment names it: "the host values were right and the UPLOAD did not land on the tensor this call reads". An opt-in device readback (`VT_DFLASH_BOUNDS_DEVICE=1`, off by default because the read synchronizes and this call sits on the route's no-sync path) now asserts the DEVICE `seq_ext` and slot-map endpoints against the host derivation | bug | -| [#2309](https://github.com/mudler/vllm.cpp/issues/2309) | `ENGINE-HYBRID-PLACEMENT` | **The fp4-resident MoE refusal was lost when W3c moved every architecture onto the shared placement seam.** `RunMoeBlockPlaced` refused the arm; the refactor left that helper dead and the live `RunMoePlaced` path accepted it. Placing an fp4-resident arm uploads every expert at load and then computes on the host across the bus, so it is SLOWER than not placing — and **a token gate cannot see it**, because the tokens stay correct and only the placement is wrong. Refusal restored as a `placeable` / `unplaceable_reason` contract on the seam itself rather than in each caller, so a newly wired architecture inherits it; callers pass `layer.moe.expert_gate_fp4.empty()`. It fires only when a placement is in force (`placed_on != engine_device`), leaving an ordinary unplaced load untouched, since a guard that fired there would break every load. Proved by a COMPILING mutation: with the guard rewritten never to fire, `test_device_placement` goes red at 1 case / 2 assertions. The first mutation attempt failed to compile under `-Werror` on the unused parameters and the stale binary reported 19/19 SUCCESS, which is a passing mutant proving nothing — the mutant build's rc=0 is part of the evidence. Found while gating W3c, fixed in the same flow. Spec [`specs/hybrid-placement.md`](../specs/hybrid-placement.md) §W3d | bug | -| [#2302](https://github.com/mudler/vllm.cpp/issues/2302) | `MODEL-DSV4-DSA-COMPOSE` | **`dsv4-dsa-compose.md` said `KV-DSV4-MULTICACHE` W3 was OWED and that this row's W1 was blocked on it; W3 landed 2026-08-27 as `ca3dcda21` ([#2078](https://github.com/mudler/vllm.cpp/pull/2078), [#2068](https://github.com/mudler/vllm.cpp/issues/2068) CLOSED), and the real blocker is the ownerless W5.** Found 30 minutes after that spec merged as `a4b333329` ([#2287](https://github.com/mudler/vllm.cpp/pull/2287)), while verifying the row's readiness against the tree rather than its records. **The code names the wall itself** -- `ModelRegistry::Forward` (`src/vllm/model_executor/models/model_registry.cpp:430-440`) refuses with "no registered forward consumes a cache set keyed by layer name ... row KV-DSV4-MULTICACHE W5 owns the consuming forward" -- so a DeepSeek-V4 engine today constructs, publishes AND allocates all 167 buffers and refuses at the first forward. **This makes the ordering HARDER, not softer:** W3 had an owner and landed, while W4-W7 are proposals with no owner at all, so nothing in `MODEL-DSV4-DSA-COMPOSE` can begin until W5 acquires one. **Cause, which is the reusable part: two stale records agreed with each other and neither was the tree.** #1925's index row predates W3 and still describes the runner dropping groups silently; `kv-dsv4-multicache.md` `## Now` opened with "W3 (#2068) is claimed" while its OWN closing paragraph already said the engine allocates all 167 buffers and refuses naming W5. AGENTS.md `## History is git` says "Before you conclude anything about past work, read the spec and run `git log -S`", and `git log --oneline --grep '2068'` shows `ca3dcda21` immediately -- it was not run. FIXED by correcting the `## Dependencies` table, `## Now`, the `## Work breakdown` prerequisite and the `## Stop conditions` entry in [dsv4-dsa-compose.md](../specs/dsv4-dsa-compose.md), the `MODEL-DSV4-DSA-COMPOSE` row in [kernel-matrix.md](../kernel-matrix.md), and the misleading opening sentence in [kv-dsv4-multicache.md](../specs/kv-dsv4-multicache.md) so the next reader is not caught the same way. The [#2286](https://github.com/mudler/vllm.cpp/issues/2286) index row carries the original wrong claim and is NOT edited, because the index is append-only by policy; this row and the corrected spec are the record | bug | -| [#2307](https://github.com/mudler/vllm.cpp/issues/2307) | `SPEC-DFLASH2` | **`VT_DFLASH_BOUNDS_DEVICE` was read from `src/` and documented nowhere, so `check-env-doc` was RED on `origin/main` and every branch cut from it inherited a red preflight.** Introduced by `21ef6f053` ([#2274](https://github.com/mudler/vllm.cpp/issues/2274), [#2304](https://github.com/mudler/vllm.cpp/pull/2304)), read at `src/vllm/model_executor/models/qwen3_dflash_internal.h:375`. Reproduced on a CLEAN `origin/main` checkout with no local changes, so it was not an in-flight artifact; it failed both `check-env-doc` and `test_check_env_doc` under `scripts/agent-preflight.sh`. FOUND while landing [#2302](https://github.com/mudler/vllm.cpp/issues/2302) and FIXED IN FLOW, per AGENTS.md's rule that filing does not defer the fix. **Documented in `docs/ENVIRONMENT.md` rather than allowlisted**, following the convention its own family sets -- `VT_DFLASH_PAGED`, `VT_DFLASH_GRAPH`, `VT_DFLASH_ATTN_BLOCK` and `VT_FA2_DFLASH_BLOCK` are all documented there, and the allowlist is for kernel-internal tuning switches. The distinction is load-bearing here rather than clerical: the switch adds two `Copy` + `Synchronize` round-trips onto a path whose entire purpose is to avoid a sync, so enabling it changes the timing of the very thing `SPEC-DFLASH2` measures -- a DIAGNOSTIC run, never a speed run, and a reader has to be told that | bug | +| [#2309](https://github.com/mudler/vllm.cpp/issues/2309) | `ENGINE-HYBRID-PLACEMENT` | **The fp4-resident MoE refusal was lost when W3c moved every architecture onto the shared placement seam.** `RunMoeBlockPlaced` refused the arm; the refactor left that helper dead and the live `RunMoePlaced` path accepted it. Placing an fp4-resident arm uploads every expert at load and then computes on the host across the bus, so it is SLOWER than not placing — and **a token gate cannot see it**, because the tokens stay correct and only the placement is wrong. Refusal restored as a `placeable` / `unplaceable_reason` contract on the seam itself rather than in each caller, so a newly wired architecture inherits it; callers pass `layer.moe.expert_gate_fp4.empty()`. It fires only when a placement is in force (`placed_on != engine_device`), leaving an ordinary unplaced load untouched, since a guard that fired there would break every load. Proved by a COMPILING mutation: with the guard rewritten never to fire, `test_device_placement` goes red at 1 case / 2 assertions. The first mutation attempt failed to compile under `-Werror` on the unused parameters and the stale binary reported 19/19 SUCCESS, which is a passing mutant proving nothing — the mutant build's rc=0 is part of the evidence. Found while gating W3c, fixed in the same flow. Spec [`specs/hybrid-placement.md`](specs/hybrid-placement.md) §W3d | bug | | [#2312](https://github.com/mudler/vllm.cpp/issues/2312) | `SPEC-DFLASH2` | **`check-env-doc` was RED on `main`: `21ef6f053` (#2274 / #2304) landed `VT_DFLASH_BOUNDS_DEVICE` documented in its index row and its code comment but NOT in `docs/ENVIRONMENT.md`.** A BASE failure rather than a branch one — every branch cut after that commit inherits a red `scripts/agent-preflight.sh`, cannot reach a green gate before push, and the red is charged to whichever unrelated change runs the gate next; found exactly that way while gating [#2309](https://github.com/mudler/vllm.cpp/issues/2309). Documented beside the other `VT_DFLASH_*` entries as user-facing rather than allowlisted as kernel-internal, because the readback is a `Download` that SYNCHRONIZES on a path deliberately kept sync-free, so it changes timing as well as checking. Fixed in the same flow, as the in-flow rule requires | bug | -| [#2275](https://github.com/mudler/vllm.cpp/issues/2275) | `MODEL-MM-QWEN4-EXP` | **`LoadStackedExperts` implements ONE of the three residencies `GgufLoadPolicy::Route` can return for a stacked expert tensor, and silently expands the other two to bf16.** `Route` answers `kKeepQuant`, `kKeepF16` or `kNvfp4Fp4` for `GgufTensorRole::kStackedExpertWeight`, and the f16 arm is genuinely reachable for a rank-3 tower because `KeepF16KDim` gives that role a K dim (`src/vllm/model_executor/model_loader/gguf_keep_quant.cpp:59-60`). `LoadStackedExperts` (`src/vllm/model_executor/models/qwen4_exp_weights.cpp:148-167`) branches on `kKeepQuant` alone; both other values fall off the end into `return ExpandBf16(g, name, {e, n, k}, /*nk=*/true)` at `:166`. The loader therefore materialises a residency the policy did not ask for, and nothing refuses. **Loud today, and still a row.** At the released 512 x 640 x 2560 geometry the expansion is ~240 GB across the stack, so on every device this project owns it aborts as an allocation rather than answering wrongly — which is the good case. Two reasons it is still worth a number: the abort NAMES THE WRONG THING (an operator reads out-of-memory, not "the loader ignored your quantization policy", so the diagnosis leads away from the cause), and the POLARITY IS WRONG EVEN WHEN IT FITS — on a reduced expert count, a test fixture or a future smaller checkpoint the fall-through succeeds and hands back a bf16 tower where the policy asked for `kKeepF16` or NVFP4, and AGENTS.md §"Inherit vLLM defaults" records that a token gate cannot see a dtype that is too wide. **What closes it:** refuse by name any residency `LoadStackedExperts` does not implement, naming the requested residency and the tensor, which is the pattern AGENTS.md already requires of an unimplemented arm; supporting `kKeepF16` or `kNvfp4Fp4` for stacked experts is a separate implementation with its own red-first gate, and the refusal comes first so the gap is visible instead of discovered as an allocation failure. Found by W5d-4 ([#2249](https://github.com/mudler/vllm.cpp/issues/2249) item 4) while writing the MoE weight adapter and deliberately NOT fixed in that flow: it belongs to the W5a loader, not to an adapter branch, and a residency refusal needs its own red-before test. Recorded under `## Owed` in [`specs/qwen4-exp-flash-next.md`](../specs/qwen4-exp-flash-next.md) beside the NVFP4 debt, under campaign issue [#1978](https://github.com/mudler/vllm.cpp/issues/1978) | bug | +| [#2324](https://github.com/mudler/vllm.cpp/issues/2324) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **W5b-1 — the `Glm5NextTextAttention` block and the `OwnedTensor` -> host f32 bridge.** Split out of [#2241](https://github.com/mudler/vllm.cpp/issues/2241), which stays OPEN for W5b-2, because the two halves answer to different oracles: the block and the bridge answer to `transformers` v5.16.1 (`modeling_glm5_next.py:1064-1257`, sha256 `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b`) and the llama.cpp [#27752](https://github.com/ggml-org/llama.cpp/pull/27752) container and need no cache over them, while the decoder layer and the forward answer additionally to `MakeGlm5NextKVCache` and the `[T, hc_mult, hidden]` manifold. Three defects a fluent wrong port produces, each with its own discriminating case: (1) the converter SPLITS `kv_b_proj` and transposes only the K half, so K contracts over its first inner axis and V over its second — at the published geometry a swap is a shape error, so the gate also carries a SQUARE case where the untransposed reading is shape-valid and merely wrong, separating by 2.9469 over all 900 values; (2) CROSS-LAYER top-k sharing — a `shared` layer reuses the previous full layer's selection, and a layer that recomputes RUNS and emits plausible tokens, so the gate carries BOTH the correct output and a recomputing port's, both from the same oracle run, and asserts ours is the first (320 of 800 values differ, max separation 1.52); (3) the all-masked padded row is filled with `finfo.min` and NOT `-inf`, so its softmax is uniform and its output finite — the `-inf` mutation reds 49 of 160 assertions. **The bridge ANSWERS O22's open residency question**: decode ONE DSA layer at a time, 0.4654 GiB, never the tower, whose expanded form is 426.72 GiB against a ~119.63 GiB box; the 1 GiB per-tensor ceiling sits EXACTLY 4x above the largest legitimate tensor and EXACTLY 9x below the smallest expert bank, is checked from the SHAPE before any allocation, and cannot make O19 / [#2260](https://github.com/mudler/vllm.cpp/issues/2260)'s `MoeGateUpSwiGLUGroupedCuda` throw reachable because no overload takes an expert bank. RED captured first from the plausible wrong port (9/14 cases, 63/150 assertions); green 14/14 + 160 and 13/13 + 96; 22 of 24 negative mutations kill their gate, one is recorded as EQUIVALENT with its reason, and the other is the `BlockToFloat`-null arm no input can reach in this build, whose PREMISE gate is armed instead and proved by removing the Q8_0 decoder. **NOT REACHED from a production entry point** — the staged-slice disclosure is spec `## Owed` O25 and the wiring is W5b-2's | feature | | [#2317](https://github.com/mudler/vllm.cpp/issues/2317) | `ENG-RECORD-CONFLICT-SURFACES` | **GitHub does not apply the `merge=union` driver, so `.agents/issue-index.md` is a repo-wide lock and the PR it blocks gets ZERO check-runs rather than a red.** Measured 2026-08-29 at `origin/main` `df024dce4` while unblocking [#2303](https://github.com/mudler/vllm.cpp/pull/2303). The mechanism is now PROVEN rather than hypothesised, on one case's three real blobs with one variable: `git merge-tree --write-tree 9d672e3b3 df024dce4` exits 0 because it honours `.gitattributes:7`, while `git merge-file -p ours base theirs` over the identical inputs exits 1 with a conflict hunk because it does not, and GitHub's mergeability computation is on the `merge-file` side. Neither side of that case edits or deletes a row (ours appends 1, `#2300`; theirs appends 6, `#2223 #2230 #2286 #2274 #2309 #2312`), so both obey the append-only rule perfectly and conflict anyway. **19 of 23 open PRs touch this file.** GitHub has computed mergeability for 9 of the 19; **8 are unmergeable, 7 of the 8 conflict on the index, and for 6 of those the index is the ONLY conflicting path** (#2311 #2310 #2303 #2301 #2267 #2248; #2281 adds one spec file; #2168 alone is unmergeable for an unrelated surface, `kernel-matrix.md`). The consequence is SILENT, which is the part [#883](https://github.com/mudler/vllm.cpp/issues/883) does not carry: an unmergeable PR gets no check-run at all, because GitHub schedules `pull_request` workflows against `refs/pull/N/merge` and stops recomputing that ref once the merge fails. On #2303 head `9d672e3b3` the check-run `total_count` is **0**, while `refs/pull/2303/merge` still EXISTS and resolves to `7b84f5cb2`, frozen at parents `a4b333329` (a pre-conflict `main`) and `bfb4f87f3` (the FIRST of the branch's three commits), so a reader sees a merge ref plus two runs stuck at `queued`/`in_progress` on that stale head since 20:25Z and concludes the queue is slow. Nothing will ever arrive. #2248 is in the same state. This contradicts `AGENTS.md:74` ("carries `merge=union`, so two branches that each append a row merge without a conflict"), the same sentence in this file's own preamble, and the same sentence in `scripts/check-issue-index-append-only.py:4-5`; the preamble copy is frozen as `INDEX_PREAMBLE` at `scripts/check-agent-record.py:1973`, so correcting it is itself a gated change. Under `AGENTS.md` §Records the second admitted shape ("a genuinely append-only file that can union-merge") does not exist on this forge, and the surface degrades into the lock the same section forbids ("If N concurrent pull requests edit file F, that file is a lock") with N = 19. NOT FIXED and deliberately: moving the index to a per-row surface changes `check-agent-record.py` and `check-issue-index-append-only.py` semantics and needs its own spec, a red-before test and a fresh review. Three options are sketched in the issue and none is chosen: per-row files under a directory read by glob (the shape `owed_issues()` already uses); a derived-at-read-time index; or keeping the file and removing the SILENCE rather than the lock. Owned by `ENG-RECORD-CONFLICT-SURFACES` ([#364](https://github.com/mudler/vllm.cpp/issues/364)), whose spec `retire-shared-record-surfaces.md` was measured at `d928e2c3` before this file had its present shape and lists it under neither in-scope nor out-of-scope | bug | | [#883](https://github.com/mudler/vllm.cpp/issues/883) | `ENG-RECORD-CONFLICT-SURFACES` | **GitHub reports `CONFLICTING` on `.agents/issue-index.md` while local git merges it cleanly, so the union driver #846 armed does not stop forge conflicts.** Filed 2026 from the LTX-2.5 landing campaign against PR #880 and never indexed here; its row is appended now, in the same commit as [#2317](https://github.com/mudler/vllm.cpp/issues/2317), because the issue that first observed this class was itself untracked by the surface it is about, a row-key scan for it having returned zero. It measured both directions on one case, `git merge-tree --write-tree` rc=0 with zero CONFLICT lines and `git merge --no-commit --no-ff` rc=0 with the path reported modified rather than unmerged, and its operator consequence stands unchanged: a `CONFLICTING` verdict from the forge is not evidence of a conflict, so reproduce it with a local `git merge` before acting on it. It deliberately left the mechanism unestablished ("the leading hypothesis is that GitHub computes mergeability without applying `.gitattributes` merge drivers. I did not verify that") and proposed a two-throwaway-branch experiment to settle it; #2317 settles it instead with no throwaway PRs, by running `merge-tree` and `merge-file` over the same three blobs so attribute handling is the only variable. Related #364, #595, #846, #573 | bug | -| [#2324](https://github.com/mudler/vllm.cpp/issues/2324) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **W5b-1 — the `Glm5NextTextAttention` block and the `OwnedTensor` -> host f32 bridge.** Split out of [#2241](https://github.com/mudler/vllm.cpp/issues/2241), which stays OPEN for W5b-2, because the two halves answer to different oracles: the block and the bridge answer to `transformers` v5.16.1 (`modeling_glm5_next.py:1064-1257`, sha256 `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b`) and the llama.cpp [#27752](https://github.com/ggml-org/llama.cpp/pull/27752) container and need no cache over them, while the decoder layer and the forward answer additionally to `MakeGlm5NextKVCache` and the `[T, hc_mult, hidden]` manifold. Three defects a fluent wrong port produces, each with its own discriminating case: (1) the converter SPLITS `kv_b_proj` and transposes only the K half, so K contracts over its first inner axis and V over its second — at the published geometry a swap is a shape error, so the gate also carries a SQUARE case where the untransposed reading is shape-valid and merely wrong, separating by 2.9469 over all 900 values; (2) CROSS-LAYER top-k sharing — a `shared` layer reuses the previous full layer's selection, and a layer that recomputes RUNS and emits plausible tokens, so the gate carries BOTH the correct output and a recomputing port's, both from the same oracle run, and asserts ours is the first (320 of 800 values differ, max separation 1.52); (3) the all-masked padded row is filled with `finfo.min` and NOT `-inf`, so its softmax is uniform and its output finite — the `-inf` mutation reds 49 of 160 assertions. **The bridge ANSWERS O22's open residency question**: decode ONE DSA layer at a time, 0.4654 GiB, never the tower, whose expanded form is 426.72 GiB against a ~119.63 GiB box; the 1 GiB per-tensor ceiling sits EXACTLY 4x above the largest legitimate tensor and EXACTLY 9x below the smallest expert bank, is checked from the SHAPE before any allocation, and cannot make O19 / [#2260](https://github.com/mudler/vllm.cpp/issues/2260)'s `MoeGateUpSwiGLUGroupedCuda` throw reachable because no overload takes an expert bank. RED captured first from the plausible wrong port (9/14 cases, 63/150 assertions); green 14/14 + 160 and 13/13 + 96; 22 of 24 negative mutations kill their gate, one is recorded as EQUIVALENT with its reason, and the other is the `BlockToFloat`-null arm no input can reach in this build, whose PREMISE gate is armed instead and proved by removing the Q8_0 decoder. **NOT REACHED from a production entry point** — the staged-slice disclosure is spec `## Owed` O25 and the wiring is W5b-2's | feature | +| [#2275](https://github.com/mudler/vllm.cpp/issues/2275) | `MODEL-MM-QWEN4-EXP` | **`LoadStackedExperts` implements ONE of the three residencies `GgufLoadPolicy::Route` can return for a stacked expert tensor, and silently expands the other two to bf16.** `Route` answers `kKeepQuant`, `kKeepF16` or `kNvfp4Fp4` for `GgufTensorRole::kStackedExpertWeight`, and the f16 arm is genuinely reachable for a rank-3 tower because `KeepF16KDim` gives that role a K dim (`src/vllm/model_executor/model_loader/gguf_keep_quant.cpp:59-60`). `LoadStackedExperts` (`src/vllm/model_executor/models/qwen4_exp_weights.cpp:148-167`) branches on `kKeepQuant` alone; both other values fall off the end into `return ExpandBf16(g, name, {e, n, k}, /*nk=*/true)` at `:166`. The loader therefore materialises a residency the policy did not ask for, and nothing refuses. **Loud today, and still a row.** At the released 512 x 640 x 2560 geometry the expansion is ~240 GB across the stack, so on every device this project owns it aborts as an allocation rather than answering wrongly — which is the good case. Two reasons it is still worth a number: the abort NAMES THE WRONG THING (an operator reads out-of-memory, not "the loader ignored your quantization policy", so the diagnosis leads away from the cause), and the POLARITY IS WRONG EVEN WHEN IT FITS — on a reduced expert count, a test fixture or a future smaller checkpoint the fall-through succeeds and hands back a bf16 tower where the policy asked for `kKeepF16` or NVFP4, and AGENTS.md §"Inherit vLLM defaults" records that a token gate cannot see a dtype that is too wide. **What closes it:** refuse by name any residency `LoadStackedExperts` does not implement, naming the requested residency and the tensor, which is the pattern AGENTS.md already requires of an unimplemented arm; supporting `kKeepF16` or `kNvfp4Fp4` for stacked experts is a separate implementation with its own red-first gate, and the refusal comes first so the gap is visible instead of discovered as an allocation failure. Found by W5d-4 ([#2249](https://github.com/mudler/vllm.cpp/issues/2249) item 4) while writing the MoE weight adapter and deliberately NOT fixed in that flow: it belongs to the W5a loader, not to an adapter branch, and a residency refusal needs its own red-before test. Recorded under `## Owed` in [`specs/qwen4-exp-flash-next.md`](specs/qwen4-exp-flash-next.md) beside the NVFP4 debt, under campaign issue [#1978](https://github.com/mudler/vllm.cpp/issues/1978) | bug | | [#2327](https://github.com/mudler/vllm.cpp/issues/2327) | `PERF-QWEN35-STAGE-WEIGHTS` | **Qwen3.5's dense decode weights were ATS-retagged HOST pages, and reading them from the GPU cost 22.4% of decode throughput on GB10 — staging them as true device copies takes the 27B from 0.82x vLLM to PARITY on the measured prompt.** The target decode step is weight-bandwidth-bound (~50 GB swept per forward against GB10's 273 GB/s roof, a ~184 ms floor; activations at 8 query rows are ~100 KB), so a placement penalty on the weights IS the step time. `ResidentWeight` (`qwen3_5.cpp:1141-1170`) handed every decode GEMM a HOST pointer tagged with a device wherever `host_memory_is_device_addressable()` answers true, while vLLM's parameters are built inside the torch CUDA allocator (`model_loader/base_loader.py:52-54`) and are real device memory. MEASURED on `dgx:gpu0`, one boot, one binary at `21ef6f053`, Qwen3.8-27B bf16 + DFlash2 k=7 at concurrency 1, `VT_QWEN35_ALIAS_HOST_WEIGHTS` the only variable, four warm repeats per leg, interleaved A,B,A,B,A: alias ON medians **11.677 / 11.693 / 11.690** (the third being the TERMINAL CONTROL, matching the opening arm, so the run is admissible and drift cannot masquerade as effect) against staged **14.288 / 14.337** — **+22.4%**, with vLLM on the same prompt at **14.305** and staged ours at **14.31**. This is the mechanism `laguna.cpp:130-132` already records and already shipped default-ON for two other models (Laguna to 1.03x vLLM, DeepSeek-V4 to 1.144x); Qwen3.5 never received it, and `docs/ENVIRONMENT.md:259` names the decode risk in the alias's own row and leaves it open. **It CANNOT be a blanket default flip**, because this file also serves `Qwen3.8-2.4T-A95B` and [#1299](https://github.com/mudler/vllm.cpp/issues/1299) measured that checkpoint exhausting a 119.631 GiB box precisely BECAUSE the CUDA arm paid for its weights twice — host bytes plus a device copy. So the policy asks the BOX rather than answering once for the file: `DeviceStagingFitsBudget` stages only while `VT_QWEN35_STAGE_MIN_FREE_FRAC` (default 0.55) of total device memory remains free AFTER the copy, which a 50 GiB model on a 119.6 GiB box satisfies and the 2.4T model — already past the floor when its first dense weight arrives — never does, preserving #1299's invariant exactly. `VT_QWEN35_ALIAS_HOST_WEIGHTS=1` pins the retag, `=0` forces staging, and an unanswerable `DeviceMemoryInfo` keeps today's behaviour rather than guessing, because an unknown budget is not a licence to double a model's residency. The budget arithmetic is extracted as a PURE function so it is gateable without a device (a fake `vt::Backend` would need every pure virtual stubbed and would gate less), and five cases pin it: a fitting model stages, the #1299 shape is refused, the floor is load-bearing, an unanswerable budget refuses, and a weight larger than free memory refuses. Red-first by mutation — deleting the reserve floor fails exactly the #1299 case and the floor case, `BUILD rc=0 / TEST rc=1`. One case was WRONG when first written and is recorded as such: it asserted the EXACT floor boundary, which gates the binary representation of 0.55 rather than the policy, and was replaced by clearly-above and clearly-below cases plus a floor-moves-the-answer case. The declarations sit at the END of the header deliberately: the first draft inserted them mid-file and shifted `Fp8Weight` from `:628` to `:658`, breaking the recorded anchor three records cite and reding `check-agent-record` at 29 stale against a baseline of 28 | bug | -| [#2331](https://github.com/mudler/vllm.cpp/issues/2331) | `ENG-HYBRID-PLACEMENT` | **`main` did not build: `tests/vllm/model_executor/test_placed_moe_roundtrip.cpp` called `vllm::RunMoeBlockPlaced`, which `866075b2f` ([#2313](https://github.com/mudler/vllm.cpp/pull/2313)) deleted.** Verified at `origin/main` rather than on a branch -- the symbol was declared in ZERO files under `include/`/`src/`, the test called it 4 times, and the test was registered at `tests/CMakeLists.txt:1587`. Red on `build-test-cpu`, `build-newest-gcc` and `sanitize-cpu (thread)`, so **every branch cut from main inherited it and no open pull request could go green**, since CI builds the merge commit. #2313 was right to delete the helper -- its own message says the W3c refactor had already left it dead -- it just did not delete the test keeping it compiling. **FIXED by deleting the obsolete file, and the two cases died for DIFFERENT reasons, which is why neither was ported.** (1) The fp4-resident refusal is SUPERSEDED by a strictly stronger case `866075b2f` added in the same commit (`test_device_placement.cpp:395-440`), which reaches a REAL cross-device placement the CPU-only original could not. (2) The byte-for-byte round trip is NOT PORTABLE by construction: it worked by passing `kCPU` as the placement device explicitly, while `RunMoePlaced` reads the device from `ActiveMoePlacementPlan()` and short-circuits with `if (placed_on == engine_device) return body(engine, dh)` -- same-device placement is INERT by design. **The round-trip gate is therefore OWED, and was owed before this**: the deleted file's own header said "It does NOT prove the cross-device arm ... it is the gate W3b still owes". What the deletion removed is a file that LOOKED like coverage while exercising a helper production had stopped calling. Recorded under `## Owed` in [expert-stream-device-slots.md](../specs/expert-stream-device-slots.md) as a cross-device byte-for-byte run. Deliberately not "ported" to a same-device assertion, which would only restate `return body(engine, dh)`. Found while landing [#2302](https://github.com/mudler/vllm.cpp/issues/2302), whose docs-only pull request was red on jobs its diff cannot affect | bug | | [#2241](https://github.com/mudler/vllm.cpp/issues/2241) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **W5b-2a — the decoder layer, the mHC stream threading, `Glm5NextTextModel::Forward` and the KV binding.** The four deliverables that answer to `transformers` v5.16.1 (`modeling_glm5_next.py:1259-1329` and `:1409-1494`, sha256 `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b`) and to `MakeGlm5NextKVCache`. **The `[T, hc_mult, hidden]` manifold is what this exists to get right**: `:1477` expands the embedding to four residual streams and nothing collapses them until `hc_head` at `:1493`, and a port that threads `[T, hidden]` and collapses early RUNS — finite, right-shaped, fluent, every sublayer gate on this row still green, and no end-to-end token gate for this model exists on this fleet to catch it downstream. Gated three non-overlapping ways: the per-layer `[B, S, 4, H]` streams asserted ELEMENTWISE; the oracle's own minimum pairwise stream separation (**6.4703**) asserted so those are shown to be discriminating rather than four copies of one value; and an EARLY-COLLAPSE DECOY produced by the SAME oracle modules with the manifold collapsed to its mean and re-broadcast after every layer, which the gate asserts we differ from by the oracle's own measured **2.4032**. The fixture is a FIVE-layer mixed schedule carrying all four control-flow combinations `:1261-1272` selects between plus a `shared` DSA layer, at the PUBLISHED `hc_mult` of 4, `seq_len` 12 against `index_topk` 8, row 1 left-padded by three. **The KV binding stores the LATENT and not what the reference stores**: upstream caches the EXPANDED K/V at `:1175-1179` (32,768 values per token per layer) and `DsaCache` stores the 512-wide `k_pass` plus the 257-wide packed indexer row, which is exactly what `MakeGlm5NextKVCache`'s groups 0 and 2 publish and what a case now asserts against the production `make_kv_cache` hook on the published `config.json`; the spec decided that and said the equivalence was to be PROVED, and the proof is a case asserting `ExpandKv(a ++ b) == ExpandKv(a) ++ ExpandKv(b)` EXACTLY, which NoPE is what makes true. An 8-token prefill plus 4-token continuation reproduces the reference's own `DynamicCache` run and agrees BIT-EXACTLY with the 12-token one-shot tail. **RED first, and the red was in the ORACLE CONFIGURATION rather than the port**: 4/10 cases and 7/1647 assertions failed with layer 0 green and every DSA layer red by 2.7 to 8.1, bisected against oracle intermediates (mHC pre plus `input_layernorm` agreed to 4.8e-07, the attention did not), and the cause is that `Glm5NextPreTrainedModel` sets `_supports_sdpa = True` so a default config resolves `_attn_implementation` to `sdpa`, whose `build_attention_mask_from_topk` returns a BOOLEAN mask (`:1249-1250`) where the eager arm returns the additive `finfo.min` one (`:1252-1256`) — the two DISAGREE on a left-padded query row where every key is masked, SDPA emitting 0.0 and eager's uniform softmax the mean of the values, measured 0.0 against 0.509. The generator now pins eager, which is the arm W5b-1 gated and the only interface `:1227-1228` says a 3-D per-(query, key) mask can reach; no token gate could see this, because the rows that differ are padding. **Fourteen negative mutations, all fourteen killing their gate — after a fifteenth finding that is the one worth reading.** The mutation truncating the attention's key range under a filled cache SURVIVED at 1647/1647, because its output is all-NaN and `NaN > x` is FALSE for every x, so the running maximum in the test's `MaxGap` helper never left its initial zero and an ALL-NaN FORWARD READ AS A PERFECT MATCH on every gap assertion in the file; `MinStreamSeparation` and the cached-tail loop were blind the same way. All three now treat a non-finite value as an INFINITE gap and report the count separately, the mutation then reds 3 assertions, and the suite grew 1647 -> 1656. Green 10/10 + 1656, with the eight sibling glm5 suites unchanged and green. **NOT REACHED from a production entry point, and this wave's own scope said it would be** — `ForwardGlm5NextForConditionalGeneration` still refuses by name, so O15, O16, O17, O23 and O25 are NOT discharged: `.agents/reachability.md` is explicit that "an intermediate hop that is itself unreached does not carry", and what changed is that five separate dead ends became ONE gated assembly point. Spec `## Owed` O26 carries that in the strong form and **W5b-2b owns the wiring**, which is why #2241 stays OPEN. W5b-2b's two halves are now scoped from measurement rather than guess: the weight bridge for the KDA, MoE, dense-MLP and mHC arms, whose 42 sparse layers' routed experts are ~1,150 GiB in f32 against a ~119.63 GiB box so an on-demand per-expert decode is the only shape that fits (`kBridgeTensorF32ByteCeiling` correctly refuses a 9.0 GiB bank today, which is O25's gate working as designed); and the engine binding, which is the SMALLER half and has a house pattern — `NemotronHForCausalLM` (`nemotron_h_registry.cpp:200-213`) and `KimiLinearForCausalLM` (`kimi_linear_forward.cpp:462-476`) both carry a host arm that ignores the paged caches and re-runs the whole prefix, and a survey of every `: public LoadedModel` found NO model keeping per-request state on it | feature | | [#2329](https://github.com/mudler/vllm.cpp/issues/2329) | `PERF-QWEN35-STAGE-WEIGHTS` | **`VT_QWEN35_STAGE_MIN_FREE_FRAC` was read from `src/` and documented nowhere, so `check-env-doc` was RED on `origin/main` itself and every branch cut after `207c12932` inherited a failing preflight.** Introduced by `207c12932` ([#2327](https://github.com/mudler/vllm.cpp/issues/2327), [#2328](https://github.com/mudler/vllm.cpp/pull/2328)), read at `src/vllm/model_executor/models/qwen3_5_weights.cpp:218` and explained only in a code comment. Reproduced on `origin/main`'s own bytes, extracted with `git archive` into a clean directory with no branch involved: `check-env-doc` rc 1 naming that one variable; with the entry, rc 0 over 396 scanned names. **This is [#2312](https://github.com/mudler/vllm.cpp/issues/2312) recurring, the same class within one day** -- that row records the identical failure for `VT_DFLASH_BOUNDS_DEVICE` from [#2304](https://github.com/mudler/vllm.cpp/pull/2304), fixed by [#2313](https://github.com/mudler/vllm.cpp/pull/2313) with the same note that it is a base failure every later branch inherits. Two occurrences in a day suggests the gap is STRUCTURAL rather than an oversight: nothing forces the doc entry at the point the knob is introduced, and the gate that would catch it only runs against a base that already merged. **Documented in `docs/ENVIRONMENT.md` rather than allowlisted**, mirroring #2313 and the convention this knob's own family sets -- `VT_QWEN35_ALIAS_HOST_WEIGHTS`, whose behaviour this variable governs, is documented there, and the allowlist is for kernel-internal tuning switches. It is user-facing on its face: it decides whether a dense weight is staged as a true device copy or left aliased, which is the difference #1299 measured between a model that decodes and one that exhausts a 119.631 GiB box. Semantics read from the code rather than transcribed: stage only while `free - bytes >= frac * total`, default `0.55`, with unset, empty, unparsable, `<= 0` and `>= 1` all falling back to `0.55` | bug | | [#2261](https://github.com/mudler/vllm.cpp/issues/2261) | `MODEL-MM-QWEN4-EXP` | **The G4 llama.cpp ladder cannot run: `llama-server` at the `qwen4exp` pin reports NO KV size, so `KV_BYTES_PER_TOKEN` has to be measured on a lease.** `scripts/qwen4exp-llamacpp-ladder.sh` extracted `KV self size = N` from the server log and, finding nothing, set the term to 0 and passed. Measured on the row's own production capture — `decode-proof/llama-server.log`, the COMPLETE unfiltered server output at 1,862 bytes — there is no `KV self size`, no `llama_kv_cache:` sizing line and no allocation summary at all, with sixteen minutes between `load_model:` and `threadpool init` and nothing printed in between; `/props` carries no KV bytes either, its only sizing fields being `n_ctx = 4096` and `total_slots = 1`. `KV_BYTES_PER_TOKEN` also defaults to 0, so on the real box NEITHER check carried a KV term while the ladder configures `CTX_TOTAL=49152` over 32 slots against a 67.5 GiB model on a 119 GiB unified-memory device that reboots rather than swaps — the 128-GiB-on-a-119-GB-box family, wearing the guard's own name. Nothing caught it because every server stub in the test suite emitted the line, so the fixture and the measured denominator disagreed. The guard now REFUSES (`E_KV_UNREPORTED`, 21) when neither the engine nor the operator supplies a term, and the fixture is silent about KV as the real server is. `/metrics` was read and rejected (it publishes a usage RATIO, not a size); a post-launch RSS check was rejected because whether GB10's unified `cudaMalloc` shows in `smaps_rollup` cannot be settled without a lease. Owed: a per-token cost measured under `-np 32 -c 49152` on a leased load | bug | | [#2262](https://github.com/mudler/vllm.cpp/issues/2262) | `MODEL-MM-QWEN4-EXP` | **The llama.cpp arm's mutation sweep is not re-executable and its CUDA toolchain is asserted rather than pinned.** Two reproducibility debts narrowed rather than closed by the repair that brought `scripts/qwen4exp-llamacpp-build-cuda.sh` and `scripts/qwen4exp-llamacpp-decode-proof.sh` into the tree. (1) `docs/bench-evidence/qwen4exp-llamacpp-ladder-arm-20260829.md` records 11 mutations red with none unarmed, and the sweep DRIVER is not committed, so nine of them are a claim about a run that happened once on one machine; two are executable tests, `test_set_u_mutation_removing_one_default_goes_red` and `test_kv_mutation_restoring_the_fail_open_default_goes_red`. (2) `apt-get install -y cuda-toolkit-13-0` pins a CHANNEL, not a version, so a rerun gets whatever apt serves and, before this, nothing would have noticed; the build now carries `EXPECT_NVCC=13.0.88`, the version the evidence records, compares it against `nvcc --version` and exits 89 on a mismatch, which makes drift visible without making apt serve one version. Owed: a committed sweep driver (or the nine as tests), and a genuinely pinned toolchain — a versioned apt pin or a recorded container image | record | -| [#2336](https://github.com/mudler/vllm.cpp/issues/2336) | `MODEL-MM-QWEN4-EXP` | **The layer loop's remaining prerequisite is the PLE BLOCK and its GATE op, and the spec's `## Now` contradicted itself about that on `c0fa299b1` — two LIVE enumerations, which is #2288 in its seventh turn.** One paragraph said "NONE remain. The count is ZERO" of #2249's five prerequisites (true, and only about those five); eleven lines below, "What has no production shape yet is the PLE block, the GDN weight adapter onto `GdnLayerWeights`, the hyper-connection stream through the per-layer loop, and the loop itself". A wave dispatched to write the loop read the first and returned `NEEDS_DECISION`. Measured on `bd90b92b0`, the count moves BOTH ways. The PLE GATE (`modeling_qwen4_exp.py:1181-1182` — `gate.abs().clamp_min(1e-6).sqrt() * gate.sign()`, then `sigmoid(gate) * value.unsqueeze(-2)`) was OP-SIZED and nothing had ever named it: a `git grep` for `clamp_min`, `signed_sqrt` and `copysign` over `src/vt include/vt` returned ZERO lines and the only implementation was the host, file-local `SignedSqrtGate`. The DOT around it needs no new op (`vt::BatchedMatmul` over `[T*hc,1,H] x [T*hc,H,1]` VIEWS), and the multiply cannot reuse `vt::SigmoidGateBf16` (refuses by element count) or `vt::MulColVecF32` (per-COLUMN) because BOTH its operands broadcast — so ONE fused op is owed, not five. Two listed items are smaller than "missing production shape" implies: the GDN adapter is a nine-assignment FIELD COPY (with `output_gate_type` sigmoid-vs-silu, a lost `ResidentWeight::d_dev`, and an empty `in_proj_ba` as its real risks) and the hyper-connection widen is `vt::IndexSelect`. What is left is the PLE BLOCK — the LAST block seam, `PleForward` has zero cross-TU callers and `vt::RmsNormGroup` has zero production callers — and the loop. Split W5e-1 (the gate op), W5e-2 (the block), W5f (the loop). **W5e-1 LANDED: `vt::Qwen4ExpPleGate`, CPU arm, gated against section J of `qwen4_exp_ple_goldens.inc`, UNREACHED by design and recorded under `## Owed`; W5e-2 and W5f remain open on this issue.** Spec [qwen4-exp-flash-next.md](../specs/qwen4-exp-flash-next.md) | bug | -| [#2337](https://github.com/mudler/vllm.cpp/issues/2337) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **W5b-2b — the weight bridge for the other four arms and the engine binding: `ModelRegistry::Forward` REACHES GLM-5.3-Flash.** Split out of [#2241](https://github.com/mudler/vllm.cpp/issues/2241) for the same reason W5b-1 was split into [#2324](https://github.com/mudler/vllm.cpp/issues/2324): #2241's index row is spent on W5b-2a and this file is append-only with one row per issue. `ForwardGlm5NextForConditionalGeneration` stops refusing by name, which discharges the reachability halves of O15, O16, O17, O23, O25 and O26 — the six debts this row has carried since W2 saying the KDA arm, the mHC bricks, the DSA indexer, the MoE block, W5b-1's attention plus bridge and W5b-2a's decoder layer were each gated and none reached from a production entry point. **The MoE half was a RESIDENCY problem and not plumbing**: one sparse layer's three expert banks are 27.0 GiB in f32 and the 42 sparse layers together are 1,134 GiB against ~119.63 GiB usable, 9.5x over, and `kBridgeTensorF32ByteCeiling` already refused one 9.0 GiB bank by name — which is that gate working, not an obstacle. It did NOT move. `num_experts_per_tok` is 8 of 288, so `MoeLayerWeights` grows a borrowed `ExpertSource*` and `DecodeOwnedTensorRowsToF32` decodes a contiguous leading-axis ROW RANGE out of a block-resident tensor: one expert is 100,663,296 f32 bytes (0.09375 GiB) and one bank row is 33,554,432 — **32x UNDER the same unchanged ceiling the bank is 9x over** — and the RANGE is checked against that ceiling, so asking for all 288 rows is refused by exactly the arithmetic that refuses the whole tensor. A block row that is not a whole number of blocks is refused by name, because a mid-block slice does not fail: the decoder reads the next block's scale and returns plausible values from the wrong quantization. `MoeForward` now visits each HIT expert ONCE, grouped, which is upstream's own order (`Glm5NextTextExperts.forward` loops the hit experts, not the tokens) and what bounds the peak at one expert; every `[t, j]` slot is still computed independently, so the resident path is byte-identical and is asserted EXACTLY against a bank-resident reference. The same row range serves the two other tensors no device here holds in float — `token_embd.weight` and `output.weight`, 2.36 GiB each — as a per-token gather and a 64 MiB-chunked head. **The binding** is `glm5_next_forward.{h,cpp}`: `Glm5NextGgufLayerSource` holds ONE layer slot and drops the previous layer before bridging the next, and `TextModelForward` grows an overload over that source so the manifold, the `prev_topk` threading and the `hc_head` collapse stay in one loop rather than being copied. The hook follows the surveyed house pattern — `NemotronHForCausalLM` and `KimiLinearForCausalLM` both ignore the paged caches and re-run the whole prefix, and no `LoadedModel` in this tree keeps per-request state — with ONE divergence in the safe direction: both precedents take `token_ids` as one sequence whatever `num_reqs` says, which attends across the request boundary, so a multi-request step is REFUSED BY NAME and ragged batching is owed. A non-CPU queue is refused too, because every primitive on this row is host f32 and `vt::MoeRouterTopK` dispatches on the queue's device. **RED FIRST from an EXISTING gate**: `test_glm5_next_scaffold`'s "the forward REFUSES BY NAME" case went red at 8 assertions the moment the hook ran, and the pin MOVED with the change the way W3's `MlaBlockDims` pin moved rather than being deleted by it. **Thirteen negative mutations on one tree, ALL thirteen killing their gate — after TWO SURVIVED the first suite and were repaired in the same branch.** The reachability mutation is the deliverable: deleting the `Glm5NextHostForward` call in the registry hook reds `test_glm5_next_forward` at 11 of 118. M12, swapping the two mHC sites in the layer source, left every logit BIT-IDENTICAL because the fixture's mHC ramps saturate every sigmoid and the Sinkhorn projection converges identically from either site — a gate that could only see it through the logits is a mute switch at that geometry, so the mapping is now asserted STRUCTURALLY with the two tensors asserted to DIFFER, and it reds 24. M13, removing the per-expert grouping, survived because the case ran ONE token, where every selected expert is hit once whatever the code does; a six-token case now fills 12 slots from 2 distinct experts and it reds 3. M5 kills by SIGSEGV (rc=139) rather than by an assertion, which corrected the refusal's own message — without it the loop dereferences a null source, it does not read zeros. Green: forward 9/118, bridge 19/32228, scaffold 38/2652, six sibling suites unchanged and green. **No token, load or speed number is claimed for the 321.32B model and none was gated** — CI runs a synthetic 4-layer miniature at `hidden_size` 32, O1 stands, and O27 names what is still owed: the `shared` indexer arm, W5b-2a's `LayerCache` binding (the full-prefix recompute does not call it), ragged batching and the device arm. Campaign [#1998](https://github.com/mudler/vllm.cpp/issues/1998) | feature | -| [#2343](https://github.com/mudler/vllm.cpp/issues/2343) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **GLM-5.3-Flash LOADS on `dgx:gpu0` and the ENGINE refuses ABOVE the model's own forward, so W5b-2b's product claim of generation is FALSE and is corrected here.** Measured 2026-08-30 by driving the production C ABI at the staged 101.2535 GiB `unsloth/GLM-5.3-Flash-GGUF UD-Q2_K_XL` artifact — `vllm-cli --device cpu --max-tokens 8` on a `vllm-cli` built at `349df8e9a` in the leased container. **The load SUCCEEDED**, which is the FIRST materialized load of this model that has ever happened: all four shards open, the tower materializes, and the engine sizes its caches — `max_model_len` auto-fits from 1048576 to 8192 against 256 blocks of 32 tokens, and `max_num_seqs` drops from 32 to 1 because one 4,390,912-byte GDN state fills a unified page of 4288 tokens. **The FIRST step then threw**, at the `input.multi_kv` guard at the TOP of `ModelRegistry::Forward`: `22 KV cache(s) from 2 published group(s) reached this forward, first 'model.layers.3.self_attn.attn', with block tables gathered for 3 of 3 published group(s), and no registered forward consumes a cache set keyed by layer name`; `vllm-cli: completion failed (status 3)`. **NO TOKEN WAS GENERATED.** That guard is KV-DSV4-MULTICACHE W3's ([#2068](https://github.com/mudler/vllm.cpp/issues/2068)) and fires for ANY model publishing a multi-cache topology BEFORE dispatch to its hook; GLM-5.3-Flash publishes three groups (W5, [#2223](https://github.com/mudler/vllm.cpp/issues/2223)), so the engine stops there and the consuming forward is that row's to write, not this one's. **What it invalidates:** W5b-2b ([#2337](https://github.com/mudler/vllm.cpp/issues/2337)) landed "LOADS AND FORWARDS on `--device cpu`" in `docs/FEATURES.md` and "A `glm5next` file LOADS and FORWARDS" in `docs/USAGE.md`, both written from the focused gate — which is exactly the reading this measurement corrects — and both are repaired in the same change, because a record correction that leaves the lie in product output is not a correction. O27's reachability discharge STANDS in the letter it was made in (the entry point dispatches to the hook when `multi_kv` is null, and deleting the call site still reds the gate at 11 of 118) and must NOT be read as "a user can generate text". **What is newly true:** a materialized load exists, correcting the first clause of O7 and of the `docs/USAGE.md` weights row; load plus engine init took under 26 minutes wall, which is a DURATION and not a throughput number. **Peak RSS is still owed and was NOT sampled** — a defect in the staging script, not a property of the run. Two further staging defects are recorded rather than hidden, because each produced a job that looked like a product result: `/usr/bin/time` is absent in the leased container (rc=127, the build never started) and a reused `/tmp/b` held a cmake cache keyed to the previous attempt's source path. Spec `## Owed` O28, campaign [#1998](https://github.com/mudler/vllm.cpp/issues/1998) | bug | -| [#2345](https://github.com/mudler/vllm.cpp/issues/2345) | `ENG-HYBRID-PLACEMENT` | **`hybrid-placement.md` claimed `RunMoeBlockPlaced` executed under `test_placed_moe_roundtrip` "byte-identical to the direct call and mutation-proven", after both had been deleted -- so the spec advertised a byte-for-byte placement gate that does not exist.** `866075b2f` ([#2309](https://github.com/mudler/vllm.cpp/issues/2309)) deleted the helper once W3c had made it dead; `6416aab85` ([#2331](https://github.com/mudler/vllm.cpp/issues/2331)) deleted the test, which was calling a symbol that no longer existed and had stopped `main` building. **The claim cannot simply be rewritten, for a structural reason:** `RunMoePlaced` short-circuits when the placement device equals the engine device (`if (placed_on == engine_device) return body(engine, dh);` -- no copy, no allocation), so the transfer path is reachable ONLY cross-device. The deleted test reached it by passing `kCPU` as the placement device explicitly, which the seam no longer accepts. FIXED by correcting the bullet to say there is no such gate and pointing at the cross-device gate recorded under `## Owed` in [expert-stream-device-slots.md](../specs/expert-stream-device-slots.md). **Filed rather than quietly edited, and this is the point of the row:** it is the fallout of #2331, which was my own change, and a record that OVERSTATES coverage is precisely the defect that cost real time hours earlier the same day -- [#2302](https://github.com/mudler/vllm.cpp/issues/2302), a wrong dependency written into a spec because two stale records agreed with each other and neither was the tree. Same shape, same treatment: a traceable correction rather than a silent one. The narrative at `hybrid-placement.md:442` is accurate HISTORY of how the code got here and is deliberately untouched; only the live-coverage claim was wrong | bug | -| [#2282](https://github.com/mudler/vllm.cpp/issues/2282) | `BACKEND-TENSTORRENT-QWEN35` | **The residency state manufactures the staging writes it then pays for: 7-8 full-tensor mesh-CQ writes per decode step, each charged the per-op CQ tax W5 measured.** The W6 probe (#2273) counted 30 persistent-route restages on a 3-token eager leg — `[11,6144]`×17 (one stable activation-hidden slot) + `[176,128]`×13 (three rotating pool bases) — and the causal chain is entirely ours: `MarkHostWritten` (`tenstorrent_ops.cpp:5654`, callers `tenstorrent_backend.cpp:56,66,70`) marks a slot host-current/device-stale, including `OnScratchBlockAcquired` where a retained DevicePool block becomes a NEW tensor whose device bytes are garbage; `CommitHost` (`:1231`) drops the device shadow entirely on a host in-place write; the next device use re-uploads the FULL tensor through the persistent arm (`:561`). The lever: eliminate staging writes instead of amortizing them — a residency state precise enough that a step stages each slot's bytes once, or zero times when the consumer overwrites the full buffer on device (candidate mechanisms: device-will-overwrite reservation, narrowed `CommitHost`, upload-on-write; the implementer derives and records the actual mechanism with explicit restage semantics per the W5 aliasing awareness). Invariant: staging stays bit-identical — sacred golden pair 16/16 STRICT, full TT suite green, speed never tokens; capture-unsafe refusals keep semantics; f32 arms keep declared dtypes; #1486 never-destroy holds; the production decode path's write count must move, observably. Evidence owed: same-method before/after on the P150 (identical leg, JIT-discard, one lock hold) reporting BOTH per-step write count and wall time, plus a fresh benchmark-record entry; a count that does not drop or a wall that does not move is a reported result — the attribution shifts or the lever is named unreachable with the trace that proves it. Owned by `BACKEND-TENSTORRENT-QWEN35`, successor to W6 (#2273, closed as inexpressible via #2280) | perf | -| [#2294](https://github.com/mudler/vllm.cpp/issues/2294) | `BACKEND-TENSTORRENT-QWEN35` | **`CopyDeviceDeviceIfCapture` records src's geometry in dst's slot without updating `dev_rows`/`dev_cols` — latent until a differing-geometry D2D copy exists.** Found by the W7 fresh review ([#2282](https://github.com/mudler/vllm.cpp/issues/2282)), audited statically at the repair head: the arm (`tenstorrent_ops.cpp:5948`) replaces dst's device shadow with a clone of src's device tensor — src's logical shape — without updating the slot's recorded geometry, and its guard only requires equal slot byte sizes (dtype-blind), so two tracked slots of equal byte size but different `[rows,cols]` geometry (or different element sizes) would make a later exact-shape `EnsureDevice2D(dst, recorded_rows, recorded_cols)` hit return src-shaped bytes to a declared-geometry consumer. No such caller exists today, so the defect is latent: every `qwen3_5.cpp` copy site (569, 1185, 1230, 1298, 1308, 1380-1389, 1467) passes a host source and is refused at `FindSlot(src) == nullptr`; the only tracked device→device callers are GatherRows-style row gathers (`qwen3.cpp:216`, `commandr.cpp:176`, `deepseek_v2.cpp:642`, the gemma/commandr family) whose destination is allocated with the source's dtype and row geometry, so slot byte equality implies equal logical geometry. Fix when it goes live: the same one-liner W7 applied to `CopyDeviceDeviceIfResident` — set `dev_rows`/`dev_cols` from `src_dev.logical_shape()` in the second lock scope — with a bit-identical audit or a test proving the differing-geometry case, since the arm runs under capture where staging semantics are strictest. Listed in the row spec's `## Owed` | bug | -| [#2353](https://github.com/mudler/vllm.cpp/issues/2353) | `KV-DSV4-MULTICACHE` | **`ModelRegistry::Forward`'s multi-cache refusal named ONE owner where THREE architectures now arrive, and never named the arriving one.** The string ended `(row KV-DSV4-MULTICACHE W5 owns the consuming forward; #1925, #2068)`, which was true by construction when W3 wrote it because DeepSeek-V4 was the only thing that could publish a multi-cache topology. Three reach it now: `DeepseekV4ForCausalLM` (7 groups, all attention; W5 does own it), `Qwen4ExpForConditionalGeneration` (3 groups, owned by `MODEL-MM-QWEN4-EXP`; `Qwen4ExpTextModel::Forward` does not exist) and `Glm5NextForConditionalGeneration` (3 groups, owned by that model's own row). That row's W5 is scoped in [kv-dsv4-multicache.md](../specs/kv-dsv4-multicache.md) `## Work breakdown` as the DeepSeek-V4 DSA-sparse path that removes `deepseek_v4.cpp`'s `(void)attn_kv`, so the clause was FALSE for two of the three. **The cost is measured**: [#2343](https://github.com/mudler/vllm.cpp/issues/2343) drove GLM-5.3-Flash on `dgx:gpu0` on 2026-08-30 and stopped at this guard, and its index row, `docs/FEATURES.md`, `docs/USAGE.md` and [CLAIM-GLM53-FLASH-W5B2B.md](../claims/CLAIM-GLM53-FLASH-W5B2B.md) then each had to reconstruct in prose what the string should have said. **The repair names the architecture and computes it** from `model.registration().architecture`, the handle this function already holds, and does NOT enumerate the three rows: a hard-coded list in a refusal is exactly the construct [#2288](https://github.com/mudler/vllm.cpp/issues/2288) has driven stale six times on the sibling row, in both polarities. Two stale anchors are repaired with it -- the comment cited `deepseek_v4.cpp:2886-2887,:2959-2960` for the discarded `attn_kv` and the values are `:3033-3034,:3105-3106`, 147 lines on -- and the ownership pin in `tests/vllm/v1/worker/test_runner.cpp` MOVED from `KV-DSV4-MULTICACHE W5` to `W3` rather than being deleted, with a negative assertion that the old clause is gone. **The guard is NOT lifted and the reason is measured**: all three arriving forwards would discard the caches (`deepseek_v4.cpp:3033-3034,:3105-3106` and `glm5_next_registry.cpp:156-157` are literal `(void)`, `qwen4_exp_registry.cpp:142` refuses unconditionally), so lifting it trades a refusal for a silent full-prefix recompute. Owned by `KV-DSV4-MULTICACHE`, FIXED IN FLOW and closed by this change; the lift and the by-name channel's recurrent gap are carried under that spec's `## Owed` | bug | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 05c406f5f9..7bad142345 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -57,14 +57,14 @@ forensics: roadmap_v1.md and the parity ledger. | `ENG-SCHED-CORE` | Text-generation running-first unified scheduler, FCFS, token budget, output update; two budget-fitting prefills co-schedule into one step (c2 parity — mirrors upstream, no divergence) | T0 | `vllm/v1/core/sched/scheduler.py:396,640,1501` @ `e24d1b24`/`702f481`; `tests/v1/core/test_scheduler.py:86,847` | `src/vllm/v1/core/sched/scheduler.cpp:114,234,365` | `tests/vllm/v1/test_scheduler.cpp:143,205,241,416`; `tests/vllm/v1/test_engine_core.cpp:271` | `planned: specs/unified-scheduler.md`; verdict [scheduler-prefill-coschedule.md](specs/scheduler-prefill-coschedule.md) | `ANCHOR-BACKFILL` | - | | `ENG-CHUNKED-PREFILL` | Basic token-budget chunked prefill | T0 | `vllm/config/scheduler.py:84`; `vllm/v1/core/sched/scheduler.py:835`; `tests/v1/core/test_scheduler.py:185,503,903` | `src/vllm/v1/core/sched/scheduler.cpp:225,548` | `tests/vllm/v1/test_scheduler.cpp:192`; `tests/vllm/models/test_qwen27_paged_forward.cpp:492` | `planned: specs/chunked-prefill.md` | `ANCHOR-BACKFILL` | - | | `KV-PREFIX-CACHE` | APC hashes, lookup, allocation, partial blocks, eviction, plus explicit/model-default cache policy. W0 ports arbitrary-group no-prefix coordination and makes hybrid/attention-free defaults cache-off. **Full-surface re-audit 2026-07-22 ([spike](specs/prefix-prompt-caching-parity.md)) — the ported core is DEEPER than this row read (chain hashing, pool, all three coordinators, the complete hybrid intersection and four single-type managers), and the residual gaps are narrower and DIFFERENT:** **`generate_block_hash_extra_keys`: W2 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC`)** — the hardcoded no-op is replaced by a 1:1 port of `kv_cache_utils.py:451-591` (`_gen_mm_extra_hash_keys` + LoRA name + `cache_salt`, fixed order lora->mm->salt; prompt_embeds deferred, no prompt-embeds path). `Request`/`EngineCoreRequest` carry `cache_salt` + `lora_name`; `FromEngineCoreRequest` sets them BEFORE the first hash (fixed a latent ordering bug: mm_features were assigned after the ctor already hashed). The latent correctness trap is CLOSED and RED-first proven: with the stub, a tenant-B request false-hits tenant-A's 48 cached tokens (`n1==48`); with extra keys `n1==0` (no false-share). This unblocks the MM + LoRA cache consumers. **prefix-cache statistics: CLOSED 2026-07-22** (W1) — `PrefixCacheStats`/`CachingMetrics` ported 1:1 with `log_stats` DEFAULTED ON, which unblocks the `BACKEND-GATE-CUDA-SGLANG-PREFIX` hit-proof requirement; first measured hit rate 0.75 on a repeated-prefix corpus; no `cache_salt`; 1 of upstream's 4 hash algos; `skip_reading_prefix_cache` absent; partial-block primitives throw (upstream's own are DEAD CODE — no caller in `vllm/` — so they are NOT owed as live behaviour). **Also cleared: the "blocked on a supported non-hybrid family" blocker is STALE** — dense models default APC ON and five have landed, yet NO gate has ever run cache-ON **MLA prefix-cache-hit assert fixed 2026-07-23** (`CLAIM-MLA-PREFIX-CACHE-ASSERT`): `FullAttentionManager::find_longest_cache_hit` asserted `kind()==kFullAttention`, aborting DeepSeek-V2 (MLA group, kind `kMlaAttention`, APC default-ON) under asserts-enabled builds — latent since `ec6f4be`, inert under Release/NDEBUG. Relaxed to upstream's precondition `isinstance(spec, FullAttentionSpec or ChunkedLocalAttentionSpec)` (single_type_kv_cache_manager.py:578-582; MLAAttentionSpec IS-A FullAttentionSpec) ⇒ accept `kFullAttention` / `kMlaAttention` / `kChunkedLocalAttention`; restores DeepSeek-V2 SACRED gate 8/8 asserts-on, full-attention byte-identical, new MLA prefix-cache-hit unit cases. | T0 | `vllm/config/model.py:1805-1860`; `vllm/engine/arg_utils.py:510,1160-1166,2473-2508`; `vllm/config/cache.py:39,93,95`; extra keys `vllm/v1/core/kv_cache_utils.py:539-574`; hasher factory `:673-730`; `vllm/v1/core/kv_cache_coordinator.py:377-425,782-834`; `tests/v1/core/test_prefix_caching.py:225,1475,2781` | hashes/managers `src/vllm/v1/core/kv_cache_utils.cpp:259,291`; **extra_keys** `generate_block_hash_extra_keys` + `_gen_mm_extra_hash_keys` `src/vllm/v1/core/kv_cache_utils.cpp`; `cache_salt`/`lora_name` on `include/vllm/v1/request.h` + `include/vllm/v1/engine/types.h`, copied in `src/vllm/v1/request.cpp` `FromEngineCoreRequest` (fields set before the first hash); `src/vllm/v1/core/kv_cache_manager.cpp:124`; no-prefix coordinator/factory `src/vllm/v1/core/kv_cache_coordinator.cpp:260,273,279,545`; model-default/hasher selection `src/vllm/entrypoints/model_loader.cpp:109,167,180,191`; CLI `examples/server/main.cpp:126`; **statistics** `include/vllm/v1/metrics/stats.h`, recorded `src/vllm/v1/core/kv_cache_manager.cpp:139-147`, reset flag `:270-276`, take-and-swap `make_prefix_cache_stats()`, per-step window fold at the end of `Scheduler::schedule()`, accessors `Scheduler`/`EngineCore`/`LLMEngine::prefix_cache_metrics()`; `Request::num_preemptions` un-deferred (`include/vllm/v1/request.h`, incremented in `Scheduler::preempt_request`) | existing APC primitives `tests/vllm/v1/test_kv_cache_utils.cpp:411,516,536`; no-prefix hybrid allocation/no-hit `tests/vllm/v1/test_kv_cache_coordinator.cpp:213`; default/override resolution `tests/vllm/entrypoints/test_loaded_engine_dense.cpp:343`; server help and online cache-off contracts `examples/CMakeLists.txt:34`; `tests/tools/test_online_gate_client.py:582,633`; statistics plus the first MEASURED hit rate `tests/vllm/v1/test_prefix_cache_stats.cpp` 12/12; **W2 extra_keys** — ported mm/lora/salt cases + ordering + hash-level no-false-share `tests/vllm/v1/test_kv_cache_utils.cpp` (29/29), manager-level salt-partition no-false-share (RED-proven `n1 48->0`) `tests/vllm/v1/test_kv_cache_manager.cpp` (10/10), CPU gate on dgx GB10. **W3 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC-W3`, dgx GB10, NOT pushed) — the FIRST-EVER cache-ON model gate:** `tests/parity/test_qwen3_apc_e2e.cpp` on `Qwen/Qwen3-4B` (dense, full-attention, APC-default-ON) 2/2 cases, 84/84 asserts — APC-ON hits 2240/2777 (rate 0.807) / APC-OFF 0; APC-ON == APC-OFF token-exact 5/6 (1 diff a vLLM-confirmed 0.125-nat near-tie); == vLLM-APC-ON teacher-forced (OFF 6/6 gap 0.0, ON 6/6 gap ≤0.125 nats, 0 outside top-20); TTFT 70.1→39.9 ms = 1.76×. NO engine code changed (gate-only over the already-shipped default-ON path); 4B SACRED 16/16 no-regression. Oracle vLLM 0.25.0. Ledger: [parity-ledger.md#L746](parity-ledger.md#L746) | [prefix-prompt-caching-parity.md](specs/prefix-prompt-caching-parity.md) (umbrella); [prefix-caching.md](specs/prefix-caching.md) (cache-policy leaf) | `DONE` (dense APC path; W4 events/W5 partial/W6 mamba-align/W7 reset endpoint tracked in `KV-EVENTS`/`KV-MAMBA-ALIGN`/own future rows) | `a41af480` | -| `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | +| `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:640` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:580`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` | | `ENG-PREEMPT-RECOMPUTE` | FCFS tail preemption with recompute | T0 | `vllm/v1/core/sched/scheduler.py:1142`; `tests/v1/core/test_scheduler.py:930` | `src/vllm/v1/core/sched/scheduler.cpp:102,157`; `src/vllm/v1/core/sched/request_queue.cpp:36` | `tests/vllm/v1/test_scheduler.cpp:247,295`; `tests/vllm/v1/test_request_queue.cpp:91` | `planned: specs/preemption.md` | `ANCHOR-BACKFILL` | - | | `ENG-CUDAGRAPH` | Decode graph capture/replay modes (host-cluster cleanup: capture-size set derived from `max_num_seqs` mirroring vLLM `_set_cudagraph_sizes`; 2026-07-18 graph-baked-scratch use-after-free fix — the 35B c2+ online-serving IMA blocker) | T0 | `vllm/config/compilation.py:53,1319,683-684,1438-1444`; `vllm/config/vllm.py:1667-1770`; `vllm/v1/worker/gpu/cudagraph_utils.py:116`; `tests/compile/test_config.py:122,229` | `src/vt/cuda/cuda_backend.cu:76,97,105`; `include/vllm/model_executor/models/decode_graph_sizes.h`; `src/vllm/model_executor/models/qwen3_5.cpp:3754,3952`; `src/vllm/v1/worker/gpu/runner.cpp:577,597`; graph-safe scratch (retire-on-grow so graph-baked scratch pointers stay valid) `src/vt/cuda/graph_safe_scratch.h`, `src/vt/cuda/cuda_moe_marlin.cu:75`, `src/vt/cuda/cuda_matmul_nvfp4.cu:766`, `src/vt/cuda/cuda_matmul_nvfp4_cutlass.cu:105`, `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:95` | `tests/vt/test_cuda_backend.cpp:98`; `tests/vllm/models/test_decode_graph_sizes.cpp`; `tests/vt/test_graph_safe_scratch.cpp`; explicit 35B gate `tests/parity/test_qwen36_paged_engine.cpp:140` | [blocktable-host-cluster-cleanup.md](specs/blocktable-host-cluster-cleanup.md); [decode-graph-scratch-uaf-2026-07-18.md](specs/decode-graph-scratch-uaf-2026-07-18.md) | `PARTIAL` | **PREFILL capture REFUTED as a lever (2026-08-17, [#1161](https://github.com/mudler/vllm.cpp/issues/1161)).** vLLM's v1 default already captures prefill piecewise (`vllm/config/compilation.py:60-63,615,630` @ `555967922`) and it is in our denominator; SGLang reached the same coverage without `torch.compile` via BCG (`SGLANG-BCG` in [sglang-matrix.md](sglang-matrix.md)). Neither helps us: GB10 2026-07-09 measured prefill GPU-idle-between-launches at **3.8%** with GPU-busy >96% on both arms, and the 27B prefill gap at **92.5% non-GEMM glue GPU work** with the dominant GEMM at +0.17% and attention AHEAD. There are no launch bubbles in our prefill to collapse. Row stays `PARTIAL`; the real residuals are exec dedup ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) and the break-point seam ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)). Spec [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | | `ENG-CUDAGRAPH-DEDUP` | Graph-executable dedup: hash each captured graph's topology and re-point ONE `cudaGraphExec` with `cudaGraphExecUpdate` on a signature hit, instead of instantiating one exec per padded bucket per model. A memory and capture-time change, NOT a throughput change — a deduped replay launches the same nodes, and the load-bearing gate is byte-identity rather than a ratio | T2 | vLLM has no analogue (its execs come from `torch.compile`, `vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); secondary oracle SGLang `python/sglang/srt/model_executor/runner_backend/cuda_graph_dedup_mixin.py:27-37,105-179,219-242,258-275,353-358` @ `f63458b5be` ([oracles/sglang.md](oracles/sglang.md)) | W1+W2 landing here behind `VT_CUDA_GRAPH_DEDUP`, default OFF until the device A/B measures the per-switch update cost: a device-agnostic dedup registry shared by both accelerator backends plus one CUDA/HIP ops table written once, wired into `EndCaptureGraph`/`ReplayGraph`/`DestroyGraph`. Baseline it replaces: `src/vt/cuda/cuda_backend.cu:222-232` instantiates a fresh exec per capture and destroys the raw graph, over the 7 (`max_num_seqs=32`) or 11 (64) buckets of `include/vllm/model_executor/models/decode_graph_sizes.h:32-41`, times NINE drivers (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight, missing the DFlash draft graph `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106`) | `tests/vt/test_graph_dedup.cpp` 13/13 cases, 65 assertions, RED-first (written and run against an absent header, and the four cases added by the fresh review of #1178, three of them run against the unfixed source) and gated on every platform via a fake ops table whose launch log makes "the right nodes ran" an observable sequence over MORE than one replay per shape; 13/13 negative mutations detected (9 at implementation, 4 at review repair). That count covers `src/vt/graph_dedup.h` ONLY. `src/vt/graph_dedup_runtime.h` had NO executable coverage on any tier, and [#1184](https://github.com/mudler/vllm.cpp/issues/1184) is what hid in that gap: the file is DESIGNED to see runtime calls fail — a refused `cudaGraphExecUpdate` probe is the feature working — and never consumed the runtime's latched error, so the next unrelated kernel reported the refusal as its own failure and every `VT_CUDA_GRAPH_DEDUP=1` run died 6/6 on GB10 as `greedy_argmax launch: invalid device function` from a launch that had succeeded. Repaired structurally rather than at twelve sites: the clear lives in `ScopedLatchClear`'s destructor (`src/vt/graph_dedup_latch.h`) installed at the six `GraphDedupOps` entry points by `MakeLatchGuardedOps`, the table's only constructor, so no raw function address reaches a field and an unwired seventh operation leaves a null the registry refuses; one line covers CUDA and HIP. The device-free half of the signature walk moved to `src/vt/graph_dedup_signature.h` and is gated by `tests/vt/test_graph_dedup_runtime.cpp` 13/13 cases, 51 assertions, RED-first against the pre-fix guard (22 failed assertions reproducing the production message), 7/7 negative mutations detected — Kahn ordering, topological re-index, sorted edge emission, the depth-4 child bound and the four graph-level escapes. STILL compile-gated only: the five node-payload cases behind the device policy. **DEVICE A/B DELIVERED 2026-08-18 on `dgx:gpu0` (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b), and it SPLIT.** Gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f` — the row squashed, so the gated tree is not an ancestor of the merge and that sha equality is what carries the claim. CORRECTNESS PASSES: 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` in every cell log where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and the three workloads hashing to three DIFFERENT values, so the identity is not vacuous. #1184 is closed by this run, because a CPU suite drives a fake runtime and cannot observe the real latched error. THE BENEFIT IS REFUTED for the case this row was filed for: `N == M` in every ON cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry's count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact the first attempt produced. Cause pre-registered before the run and then confirmed, structural rather than a tuning miss: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension sits in the KEY, no candidate group ever forms and `cudaGraphExecUpdate` is NEVER ATTEMPTED. That contradicts this row's own premise — `graph_dedup.h`'s header says the fold is for "two padded batch sizes … the same node topology with different parameters" — and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded: clocks unpinned AND the ON arm allocated exactly as many executables as OFF. Honest gaps: per-shape replay counts are unavailable (the driver prints a TOTAL, so B's ~30-per-shape is arithmetic); the driver's "N captured size(s)" counts SLOTS not captures (A reports 6, emits 3); the container's own cuBLASLt was never re-tested at CUDA 13.0 because the staged cu130 prefix was probed first and worked; only the Qwen3 dense decode driver was exercised. STILL OWED: the default flip, now NOT JUSTIFIED on this evidence rather than merely ungated; a COARSER key that could group two decode buckets at all, which the probe-before-fold design makes a cost question rather than an obviously unsafe one ([#1226](https://github.com/mudler/vllm.cpp/issues/1226), the next traceable hypothesis, deliberately NOT decided by this record); device-tier signature stability/discrimination tests; probing `current_raw` instead of `raws.front()` to retire the update-transitivity assumption; the ROCm compile; a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`); and reaching the feature from the default serving path at all — the async runner captures no decode graph, **W5, THE SAME DAY, CONFIRMED THE HYPOTHESIS THAT NEGATIVE PRODUCED ([#1226](https://github.com/mudler/vllm.cpp/issues/1226) DELIVERED).** Same box, `rc-worker-4b8lj`, boot_id `3fd9745a-d25a-426c-ba3c-97c958a85515` at both ends, GB10, driver `580.173.02`, `### DONE_AB_KEY 2026-08-18T20:58:46Z`, binary sha256 `ca114abb…c772ad` from `b48b51df1` (tar sha256 asserted before extraction). Drop the launch dimensions and the memcpy extents from the key and every bucket folds: `a_coarse` 3 graphs to 2 execs, `b_coarse` 2 to 1, `c_coarse` 2 to 1, each `probes=1 refused=0`, against `probes=0 refused=0` in every EXACT cell. **`probes=0` in the EXACT cells is the direct process-level proof of W4's source-level diagnosis** — with the launch dimensions in the key no candidate group forms and `cudaGraphExecUpdate` is never asked; drop them and it is asked once per fold and ACCEPTED EVERY TIME. The saving W4 recorded as unreachable is reachable via the key. Byte-identity holds on A (five cells, `59ebff4a…`) and C (four cells, `ff205260…`). **Workload B is VOID rather than a pass, and its cause is a NEW DEFECT that is not this row's:** the two `VT_CUDA_GRAPH_DEDUP`-unset control cells DISAGREED (`5973c5a1…` 2638 bytes vs `4cf79230…` 2650 bytes) on one binary, one workload, greedy `--temperature 0 --seed 777` at `--concurrency 16`, 23 s apart — 672 tokens both, so the byte delta is JSON width and not a length; exactly rows 17 and 18 of 21 differ, both mid-decode, both in the ragged tail `21 % 16` leaves. B's `b_off_a == b_exact` and `b_off_a == b_coarse_a` therefore compare against a baseline that does not reproduce itself and are WORTHLESS; only the OFF/OFF control made that visible, and without it B would have read as three more confirmations. Filed [#1283](https://github.com/mudler/vllm.cpp/issues/1283). **Caveats that bound this result:** nvcc was `13.3.73` here and `13.0.88` for the W4 baseline the recorded dgx gate stack names, so the OFF-vs-ON and EXACT-vs-COARSE comparisons WITHIN this binary are valid while this run and that baseline are NOT directly comparable; clocks unpinned (2405 MHz current, 3003 max, 2418 applications) and nothing measured bytes, so NO throughput and NO memory number is claimed or implied; only the Qwen3 dense decode driver was exercised; `refused=0` is ONE driver on ONE hardware and toolkit pair, which is no more a floor than W4's negative was a ceiling; and the coarse key is behind `VT_CUDA_GRAPH_DEDUP_COARSE_KEY`, default OFF, inside a default-OFF flag, on **PR [#1232](https://github.com/mudler/vllm.cpp/pull/1232) which is STILL A DRAFT — nothing on `main` folds today.** **Row stays `ACTIVE`, argued:** not `DONE`, because the fold is unreachable on every shipping configuration and the row's stated MEMORY saving has never been measured in bytes on either key; not `PARTIAL`, because nothing upstream is omitted — the coarse key is our own extension past SGLang, which keys the fields we started from; not `BLOCKED`, because nothing external stops the next step. What is owed is now a DECISION about the default plus the byte measurement and the probe-cost-at-real-churn measurement it needs, and landing #1232 first **W6, 2026-08-19, THE DEVICE-BYTE MEASUREMENT — THE BENEFIT QUESTION IS NOW CLOSED AND THE ANSWER IS NEGATIVE.** Tested `origin/main` `2c8f53d93`, which is PR #1232 LANDED, so the "nothing on `main` folds today" caveat every earlier record carried is RETIRED and this measures a configuration that ships. Same box, `rc` job `93f783de`, pod `rc-worker-4b8lj`, boot_id `3fd9745a-…` at BOTH ends, GB10, driver `580.173.02`, nvcc **13.0.88** (the W4 baseline toolkit; W5 ran 13.3.73, so W6 and W5 are NOT directly comparable while comparisons WITHIN this one binary are valid), binary sha256 `be697268…0ce657a7`, `### DONE_BYTES 2026-08-19T04:57:19Z`, 12/12 cells exit 0, zero VOID markers. **THE FOLD ENGAGES AT THE SHIPPED BUCKET SET**, which is the churn W5 could not produce: `vllm-bench` sets `max_num_seqs = concurrency`, so W32 captured `[1 2 4 8 16 24 32]` 7-of-7 and W64 captured `[1 … 64]` 11-of-11, exactly `decode_graph_sizes.h:32-41`, against the 2-3 buckets every earlier conclusion was drawn from. COARSE folds 7 graphs to 3 execs (`probes=7 refused=3`) and 11 to 5 (`probes=22 refused=16`); EXACT folds NOTHING at `probes=0`, reproducing W4 at four times the bucket count. Token ids byte-identical across every cell of a workload INCLUDING both OFF/OFF controls (`ff0db6c6…be9d` 11720 B; `e1cbf5fc…e5d0` 57620 B) — neither workload has #1283's ragged-tail shape and neither hit it. **THE SAVING DOES NOT SURVIVE ITS OWN NULL CONTROL.** `nvidia-smi --query-compute-apps` tail median (the `--query-gpu=memory.used` axis returns `[N/A]` on this box) shows W64 IDENTICAL to the megabyte in all five cells (9737) and W32's coarse arm reading 10-23 MiB HIGHER than OFF (3252/3262 vs 3262/3275). A `cudaMemGetInfo` shim summed over every instantiate gives a nominal 13.83 MiB at 7 buckets — **0.42% of a 3.25 GiB process** — and **−0.75 MiB, i.e. NOTHING, at 11**. That nominal effect is NOT ESTABLISHED on four independent grounds: `EXACT` is a TRUE NULL (same 7 and 11 retained execs, `probes=0`, so it allocates what OFF allocates) and disagrees with OFF by 10.6-13.1 MiB against a 13.83 MiB candidate; the W64 OFF/OFF pair disagrees with ITSELF by 18.2 MiB; one instantiate recorded a NEGATIVE delta (`-5,165,056` B); and `cudaGraphExecDestroy` reclaimed `0` in EVERY cell. Per-instantiate deltas for byte-identical 404-node graphs span 0 to 10,514,432 B and 17 of 27 instantiates in one cell read exactly zero, so these are POOL-GRANULAR readings and the coarse arm's throwaway probes grow that pool exactly like retained execs do. What CAN be priced: one ~390-node executable at **2.08-4.35 MiB**, 10.0-10.6 KB per node — the figure to re-run on a deep checkpoint. **THE MECHANISM INVERTS THIS ROW'S PREMISE.** The driver refuses **43% of probes at 7 buckets and 73% at 11**, every one of them `probe refused a fold (err=910 result=2)` = `cudaErrorGraphExecUpdateFailure` / `cudaGraphExecUpdateErrorTopologyChanged`. The shim's `cudaGraphGetNodes` reading says why false candidates form: the decode graphs are **TWO topologies, 376 and 404 nodes**, mixed across the buckets (`w32_off_a` captured `404 404 376 376 404 404 404`). Every refusal is about TOPOLOGY, never a parameter, so a COARSER key produces MORE false hits rather than more folds — the opposite of what W5's 2-bucket A/B suggested, and W5's `refused=0` is now explained as an artefact of workloads whose buckets only ever SHRANK, so exactly one pair was ever presented. **COST:** W32 OFF 7 instantiates / 0 updates vs COARSE 10 (3 retained + 7 probes) / 11 updates; W64 OFF 11 / 0 vs COARSE **27** (5 retained + 22 probes) / 28 updates — **2.45x the instantiate calls** to retain 6 fewer executables. **Peak transient did NOT double** — in every ON cell live-bytes peak == end, because `Register` destroys the probe before returning, so the feared "double the peak to save the steady state" trade did not occur. **A replay-time re-point DID occur** — 4 and 6 non-probe updates over 88 and 244 replays, ARITHMETIC over two printed totals and not a counter — with every cell exiting 0 and byte-identical, so `Replay`'s transitivity assumption neither aborted nor changed a token; W5 recorded that case as untested. **CAVEATS THAT BOUND THIS RESULT:** the clock pin was **REFUSED inside the lease** (`The current user does not have permission to change clocks for GPU 0000000F:01:00.0`, `clocks_pinned=0`), so **NO time-based figure is attributable** and the instantiate-wall and update-wall figures in `bytes.log` are diagnostics quoted nowhere as a result; `result=2` is ONE driver, ONE GB10, ONE toolkit; only the Qwen3 dense decode driver was exercised, as in W4 and W5; `VT_ASYNC_RUNNER=0` throughout, so the feature is STILL unreachable on the DEFAULT serving path (#1179); and `cudaMemGetInfo` cannot separate an executable's own cost from the pool chunk that satisfied it. **VERDICT, DELIVERED AND NEGATIVE:** `VT_CUDA_GRAPH_DEDUP` stays default OFF, now on MEASUREMENT rather than on silence; `VT_CUDA_GRAPH_DEDUP_COARSE_KEY` alone is a **NO-OP, not merely unsupported** — `GraphDedupCoarseKeyEnabled()` (`src/vt/graph_dedup.h:114`) is read only by the signature builder (`src/vt/graph_dedup_runtime.h:177`), only from `Register`, only under `GraphDedupEnabled()` (`src/vt/cuda/cuda_backend.cu:237`), so with dedup off its sole observable is one stderr line; both on is unsupported. **NOT A CEILING.** Three things would change it and each is traceable: find where the 376/404 split comes from (the FA-2 split-KV grid is the first suspect — a capture that fixes the node set across buckets removes every refusal); an instrument that resolves a single 2-4 MiB executable against driver pool granularity (`cuMemGetAllocationGranularity` or a pool-statistics query); and the same measurement on a 60-80 layer checkpoint, where bytes scale with node count. **Row STAYS `ACTIVE`, argued, and the argument is now narrow.** The MEASUREMENT obligations are discharged and the DECISION is delivered, which is the `DONE` case and it is a real one. Three things stop the flip and none is a checker technicality: the feature is unreachable on the DEFAULT serving path, owned by `ENG-CUDAGRAPH-BREAK` (#1179) and the "nothing lands dead" half of this row; two items still sit under #1162 itself — the device-tier signature stability/discrimination tests and probing `group.current_raw` instead of `raws.front()` to retire the transitivity assumption; and the `DONE` record surface owes a `.agents/parity-ledger.md` entry, a closing-commit owner in place of the claim, an exact test anchor and the RELEASE of `CLAIM-ENG-CUDAGRAPH-DEDUP`, which is an operator act and which this record-only branch does not own. Not `PARTIAL` — nothing upstream is omitted. Not `BLOCKED` — nothing external stops the next step. Full evidence: [benchmark-record.md](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W6`, raw at `/mnt/nas_share/rc/dedup-bytes/` | [eng-cudagraph-dedup.md](specs/eng-cudagraph-dedup.md); analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-DEDUP` ([#1162](https://github.com/mudler/vllm.cpp/issues/1162)) | | `ENG-CUDAGRAPH-BREAK` | One shared `vt` capture seam that accepts BREAK POINTS, so a forward containing a host-dependent op is still graphed instead of falling out entirely — and so the NINE hand-rolled drivers become one (count corrected 2026-08-18, [#1179](https://github.com/mudler/vllm.cpp/issues/1179); `9bc4d7f44` recorded eight). **Coverage AND CORRECTNESS row, not a throughput row** | T1 | mirror vLLM `CUDAGraphMode.PIECEWISE` splitting at `splitting_ops` (`vllm/config/compilation.py:60-63,517,615,630` @ `555967922`); construction from SGLang BCG `python/sglang/srt/model_executor/runner_backend_utils/breakable_cuda_graph/breakable_cuda_graph.py:204-243,246-274,309-333,335-367` @ `f63458b5be` (decorator + runtime stream capture, no compiler); its unit suite `test/registered/cuda_graph/breakable/test_breakable_cuda_graph.py:30,172,230` (305 lines, 11 unit cases) is mapped case for case in the spec's `## Tests to port` | **W6 MOVED THE PREDICATE** ([#1374](https://github.com/mudler/vllm.cpp/issues/1374), 2026-08-19): `GPUModelRunner::execute_model` names the step's ACTUAL uniform query length once through `v1::GraphEligibleQueryLen` (`src/vllm/v1/worker/gpu/cudagraph_dispatch.h`, INERT with no caller since #442 and now called from production) and ships it on `ModelForwardInput::uniform_query_len`; the two Qwen3.5 registrations stop re-deriving that test in twenty duplicated lines each, and both key their slot ring on `(S, q, spec)`. [#1020](https://github.com/mudler/vllm.cpp/issues/1020) CLOSES on the pair, and the key half was a LIVE collision rather than the enabler #1020 called it: `S = spec_step ? B : PadToCaptureSize(B)` puts a 4-request spec step at 1+1 tokens and an 8-request padded decode on the same `S == 8` at the base commit. The widening is BOUNDED by `VT_SPEC_GRAPH_MAX_QLENS` (default 2), because reading the actual length multiplies the spec shape ceiling by `1 + k`. Seven of the nine drivers still read `pure_decode` and are byte-identical. **What did NOT move is 'except at the break points'**: no driver in this tree serves a prefill or a mixed batch under any predicate, so that needs a prefill capture driver nobody has written and whose benefit D5 already refutes on this hardware — a publishable negative, recorded in the spec's `## Owed` as a row-level item. The pre-W6 baseline it replaces: all-or-nothing, `src/vllm/v1/worker/gpu/runner.cpp:1338-1341` routing only `pure_decode`; drivers `qwen3_5.h:275`, `qwen3_5_dense.h:391`, `qwen3_moe.h:117`, `qwen3.h:243`, `deepseek_v2.h:324`, `voxtral.h:126`, plus `deepseek_v4.cpp`, `laguna.cpp` — and the spike found the NINTH already written, `src/vllm/model_executor/models/qwen3_dflash.cpp:771,1091`. The re-derivation is measured, not asserted: `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`, the persistent DEVICE input path) exists in ONE driver and `grep -c` returns 0 in `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`, which is why `src/vllm/model_executor/models/qwen3.cpp`'s `DenseDecodeGraphForward` DECLINES the graph outright when the async device-token mirror is live. **That decline is why this is also a CORRECTNESS row** ([#1179](https://github.com/mudler/vllm.cpp/issues/1179)): a SHIPPED model has already lost its decode graph to the duplication, on the driver's own measurement (`depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate`), and the fix its comment names is the sibling's `StepDevInputs`. The row still makes NO throughput claim: the prefill refutation on the `ENG-CUDAGRAPH` row (3.8% host idle, >96% GPU-busy, 92.5% glue) stands unchanged; **#1305 ADVANCED AND EXPLICITLY NOT CLOSED, and reading the tree found a larger defect than the issue described** (2026-08-19): `qwen3_moe_registry.cpp`, `deepseek_v2_registry.cpp` and `glm4_moe_lite_registry.cpp` never constructed a `detail::DeviceTokenIdsScope` and neither `qwen3_moe.cpp`'s nor `deepseek_v2.cpp`'s `EmbedInto` ever consulted one, so `ModelForwardInput::device_token_ids` reached NOTHING in either translation unit — the decode graph AND both eager arms embedded the host vector the runner's mirror arm deliberately leaves stale for decode rows. The three registries now publish the scope (the mechanism `qwen3.cpp`, `qwen3_5.cpp`, `mistral_registry.cpp`, `internlm2_registry.cpp` and `llama_registry.cpp` already use), and each decode-graph size slot holds a `vllm::StepTokenIds` (`include/vllm/model_executor/models/step_token_ids.h`) whose destination is a device buffer with a stable address, refreshed through `vt::PersistentStepInput` — host arm for the padded vector, DEVICE arm over the real prefix, both on the main queue so the second is ordered after the combine rather than racing it. That is `vt::PersistentStepInput::RefreshFromDevice`'s FIRST production caller, retiring the staged slice W4 landed with none, and it is the fix `qwen3.cpp`'s own decline comment names rather than a fifth private copy. `qwen3.cpp`'s decline is UNTOUCHED: W4 measured its recorded cause false and its real one is unidentified. **THE ISSUE SPLITS, and only one half settles.** The EAGER half is fixed and gated on all three registrations and deserves to close. The GRAPH half does not: the mechanism these two drivers now have is functionally what `qwen3.cpp` ALREADY HAD at `338cbbfd1^` — a registry scope, consumed by `EmbedInto`, copying the mirror's ids over the embed source OUTSIDE the capture — and W4 recorded at `qwen3.cpp:1083-1095` that the depth-2 graph-ON battery STILL FAILED with exactly that in place. A stable device address buys nothing while the embed stays outside the capture, which this change itself concedes. #1305's own settlement condition is that battery, it did not run, and the issue stays OPEN with the `ENG-CUDAGRAPH-BREAK` row as owner. | owed: bit-exactness vs eager on every migrated model over MORE than one replay, on a real GPU — **W2 did NOT meet it and says so**: no `rc` lease was obtainable in its window and a CPU harness cannot replay a captured segment, so it moves to W3 with the three drivers of the same shape (G1); the host-lifetime contract of `decode-graph-scratch-uaf-2026-07-18.md` enforced AT the seam — D1's INPUT half, making the intermediates a segment reads unavailable to the `DevicePool` free list, which becomes live only for the first PIECEWISE production capture (W4); the auxiliary-stream auto-join before every segment close (`:353-361`, spec D10), live at `src/vllm/model_executor/models/qwen3_5.cpp:6254-6255,6384` and `src/vllm/model_executor/models/laguna.cpp:2572-2576,2612` (W4, W5). **Delivered by W1** ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the reachability mutation (performed; deleting the call site reds `tests/vllm/models/test_qwen3_break_point.cpp` and leaves the unit suite green); the ported SGLang unit cases with their arithmetic chains and post-replay assertions; and the break-function OUTPUT writeback (`replay_fn`/`_copy_output` `breakable_cuda_graph.py:231-235,172-201`, spec D9), whose destination is a `vt::BreakSlot` the seam owns rather than a caller reference it cannot outlive **W6 gates** ([#1374](https://github.com/mudler/vllm.cpp/issues/1374)): G2 at THREE levels because the claim has three parts — the engine (`tests/vllm/v1/spec_decode/test_mtp_depth.cpp`, a real LoadedEngine/EngineCore/Scheduler/runner stack, asserting `clamped_spec_steps`, measured 0/0/1/2/4 at k=1/2/3/4/6), the driver (`tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp`, two spec shapes of equal S and different q getting two rings and two captures), and the arithmetic (`tests/vllm/v1/worker/gpu/test_cudagraph_dispatch.cpp`). Five detecting mutations, each reddening ONE level and leaving the others green, plus an over-fire control. A SIXTH mutation was NOT detected and forced a repair: the per-request verify conjunct is redundant on every model that reads the field (both are GDN hybrids whose prefill trips the first conjunct), so it moved into `GraphEligibleQueryLen` where a mutation reds 4 assertions, and the spec records it as unreached defence in depth. **G1 re-run on `thor:gpu0` (sm_110, driver 595.78, nvcc 13.0.88): 2066 assertions, 0 failed, 0 differing on all five migrated drivers — W6 moves no logit.** The ring key's own device case was BLOCKED by [#1380](https://github.com/mudler/vllm.cpp/issues/1380), a pre-existing `cudaMalloc` inside a capturing stream on the spec arm that W6 neither caused nor regressed. **#1380 is FIXED (2026-08-19, `thor:gpu0` sm_110):** a backtrace taken AT the failing `cudaMalloc` named the site as the GDN causal-conv output `dconv` in `GdnBlockPaged`, whose block lands in the same `DevicePool` SIZE CLASS as the retained `[S, vocab]` logits at the gate's shape, so the driver's one-block pre-grow met a measured demand of two. `DevicePool` now measures the per-class PEAK a step holds live above its own baseline and `PreGrowForCapture` makes the free list serve that profile before `BeginCapture`; both Qwen3.5 drivers record it per SLOT at their cold step. The device case drives one spec shape through BOTH ring slots into a replay against `Qwen3_5DenseModel::ForwardDevice`, with the two arms SEQUENCED rather than interleaved because an eager forward between the graph arm's steps deepens the shared free list and hides the defect (measured: interleaved passed 1240 assertions at the un-fixed head). **The architecture question #1380 asked FIRST is answered by measurement on TWO devices**: `thor:gpu0` (sm_110) and `dgx:gpu0` (GB10, capability 12.1, `-DVLLM_CPP_CUDA_ARCHITECTURES=121a`) give the SAME message and the same per-step shape at the red sha (507 assertions / 8 failed) and the SAME green after (6 cases / 3306 assertions / 0 failed, `0 differing, 3 replays`). SPEC-DSPARK W8's working GB10 capture is explained rather than contradicted: whether the size classes collide is arithmetic over the MODEL's dimensions, and the real 35B's `[S, vocab]` f32 at `vocab = 151936` shares a class with nothing the GDN block allocates. Also fixed in flow: [#1394](https://github.com/mudler/vllm.cpp/issues/1394), a block table shorter than the sequence it addresses, which the CPU paged attention read past SILENTLY -- deterministic SIGSEGV on one measured build at `main` and wrong-page attention on another. **#1305 (2026-08-19)**: `tests/vllm/models/test_moe_async_device_ids.cpp`, entered at `ModelRegistry::Forward` over a synthetic safetensors checkpoint for `Qwen3MoeForCausalLM` and `DeepseekV2ForCausalLM` — the production entry point, not the driver type. Three runs each: right host ids and no mirror as the reference, stale host ids and no mirror as the CONTROL that must differ, stale host ids with the truth reaching the model ONLY through `device_token_ids` as the gate. RED first at 2 cases / 65 assertions / 10 failed / exit 1, with 800 of 800 logit values differing over four steps on both architectures and every counter at 0; GREEN after at 65 of 65, exit 0. TWO mutations, each compiled clean and each restored by sha256: deleting the registry's scope line — the production call site — reds 4 assertions across both cases and puts all 800 values back, and swapping the seam's DEVICE arm for its HOST arm leaves the logits BIT IDENTICAL at 0 of 800 differing and reds only `device_refreshes` and `host_refreshes`, which is the arm no token gate can see. Neighbours green on the same binary: `test_qwen3_moe_decode_graph_seam` 228 of 228, `test_deepseek_v2_decode_graph_seam` 230 of 230, `test_qwen3_decode_graph_seam` 231 of 231, `test_voxtral_decode_graph_seam` 230 of 230, `test_breakable_graph` 265 of 265, `test_persistent_step_input` 66 of 66, `test_model_registry` 924 of 924, `test_qwen3_moe_forward` 504 of 504, `test_deepseek_v2_forward` 1052 of 1052. **NOT measured:** the depth-2 four-concurrent battery on a device, which needs a GPU and a real checkpoint; owed. **Found red on `main` and NOT caused here:** `test_qwen3_5_decode_graph_seam` exits 139 while its assertion line reads 135 of 135 passed ([#1390](https://github.com/mudler/vllm.cpp/issues/1390)); re-measured on this branch at exit 139 with the SAME crash case and site (`test_qwen3_5_decode_graph_seam.cpp:800`, `W6: two spec shapes of EQUAL S and different q get two graphs`) both WITH and WITHOUT this branch's working-tree changes, and its printed counts are not reproducible run to run on ONE unchanged binary — three consecutive runs of the same baseline binary gave 6 passed, 2 failed and 141 assertions, then no summary at all, then no summary at all. The exit code is the only stable observation, so no assertion count from that file carries a verdict. **THE FRESH REVIEW FOUND THE GATE ABOVE COVERED HALF OF WHAT THE CHANGE CLAIMS** and the repair widened it to 6 cases / 191 assertions / exit 0. What was ungated: the EAGER arms of both models — the half no graph refusal could have mitigated — and the THIRD registration, `glm4_moe_lite_registry.cpp`. Deleting the `TakeDeviceTokenIds` + `d.b.Copy` block from BOTH `EmbedInto` overloads left the old gate green at 2/2 and 65/65; deleting the GLM registry's two-line scope did too. The lane is now selected by the registry's OWN predicate: a case that constructs `StaticGraphCpu` gets the decode graph, a case that does not gets `ForwardDevice`, and `through_seam` asserts the `vt::PersistentStepInput` counters BOTH ways so a case cannot drift onto the other lane and stay green. Three detecting mutations, each compiled clean and each restored: the two `EmbedInto` call sites reds the 3 EAGER cases only (exit 1, 3/6); the GLM scope reds the 2 GLM cases only (exit 1, 4/6); the seam's `RefreshFromDevice` call reds the 3 GRAPH cases only (exit 1, 3/6). A fourth mutation FAILED TO BUILD under `-Wunused-parameter` and its verdict was DISCARDED rather than read as a pass. **Still owed, and not implied:** the behavioural half of the device contract — that the copy reads DEVICE memory, and that it is main-queue-ordered after the combine — is untestable on the CPU backend, where `Backend::Alloc` returns host-addressable memory and both refresh arms reduce to the same memcpy from the same address; swapping the device arm for the host arm leaves the logits BIT IDENTICAL and reds only the counters, which gate the instrument rather than the behaviour. | spec [eng-cudagraph-break.md](specs/eng-cudagraph-break.md) (W0 spike DONE 2026-08-18: the existing `vt` capture vocabulary `include/vt/backend.h:208-222` expresses a SEGMENTED capture with NO new virtual, because `EndCaptureGraph` stores nothing (`src/vt/cuda/cuda_backend.cu:225-232`); a break point is expressible with one `thread_local` capture pointer plus a free function, no compiler and no decorator); **W1 DONE 2026-08-18 ([#1192](https://github.com/mudler/vllm.cpp/issues/1192)): the seam LANDS** — `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` (`include/vt/breakable_graph.h`, `src/vt/breakable_graph.cpp`), the SGLang unit suite ported case for case (`tests/vt/test_breakable_graph.cpp`, 24 cases / 163 assertions, re-derived 2026-08-18 by `ninja test_breakable_graph && ./build/tests/test_breakable_graph`; the recorded 14/81 never re-derived at any head of this branch), and ONE break point registered at the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`). **The exit criterion W0 deliberately left open is ANSWERED on a leased GPU:** `cudaStreamEndCapture` then `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between is LEGAL under `cudaStreamCaptureModeThreadLocal` (`src/vt/cuda/cuda_backend.cu:204-206`) — `orin:gpu0` via an `rc` lease, driver 12060, 3 replays with fresh inputs, 0 mismatches, bare zero-work re-begin legal too. G2 reachability is `tests/vllm/models/test_qwen3_break_point.cpp`, which drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments (mutation: delete the call site ⇒ 1 segment ⇒ RED), and holds G4 in the same case at 500 logits / 0 differing bit for bit. STAGED SLICE, named: the scope and the container are not yet ENTERED from a production step — no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph` — and the spec's `## Owed` lists it with W2 as owner, alongside the D10 auxiliary-stream auto-join (W4/W5), G5's ROCm/Tenstorrent arms (W3) and G1 on a real GPU (W2). **The capture-failure drain is NOT among them: it landed HERE**, as behaviour (`std::uncaught_exceptions()` compared against the depth recorded at scope entry, so a break function or ordinary model code throwing mid-capture destroys the partial container instead of handing back a forward that reports `captured() == true`) and as three gated arms (tests 13a, 13b, 13c). The spec's `## Owed` strikes the item through and reads DELIVERED in W1; this cell said the opposite until 2026-08-18 because `cba969857` re-derived field 6 alone. **W2 DONE 2026-08-18 ([#1261](https://github.com/mudler/vllm.cpp/issues/1261)): `Qwen3DenseDecodeGraph` MIGRATED and the seam is ENTERED from a production step**, which retires W1's staged slice. `Qwen3DenseDecodeGraph::Step` opens a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` and replays through `BreakableGraph::Replay`; the hand-rolled `BeginCapture`/`EndCaptureGraph` pair, the raw `void*` handle, the `bool captured` flag, the `DestroyGraph` loop and the driver's own `VLLM_CPP_CUDAGRAPH` read are gone (re-derivation items 1, 2, 5, 6). The migration ADDED `vt::GraphCaptureMode`, mirroring vLLM's `CUDAGraphMode` (`vllm/config/compilation.py:59-63`), whose v1 default `FULL_AND_PIECEWISE` (`:63`) is documented at `:630-632` as a FULL graph for DECODE batches and a piecewise one for prefill/mixed, with `decode_mode()` (`:65-66`) selecting the full half and the runtime reading it at `vllm/v1/worker/gpu/cudagraph_utils.py:185-186`. A decode driver opened `kPiecewise` would have turned a fully graphed decode step into ONE EAGER ATTENTION CALL PER LAYER between graph replays — not vLLM's decode behaviour, and invisible to every token gate here. `GraphBreak` in a `kFull` scope takes the pass-through arm and `AppendBreak` REFUSES a registration in that mode. G2 is `tests/vllm/models/test_qwen3_decode_graph_seam.cpp` (3 cases / 124 assertions), which asserts the SEAM's counters because a driver calling `Backend::ReplayGraph` directly leaves an identical backend log; the mutation restoring the pre-W2 raw pair (18 lines, compiled clean) left `test_breakable_graph` 27/27, `test_qwen3_break_point` 2/2 and `test_qwen3_forward` 10/10 GREEN and reddened only this file. G4 in the same file: capture step vs `Qwen3DenseModel::Forward`, 100 logits, 0 differing. **The async decline at `qwen3.cpp` STANDS and is now GATED in both arms**: migrating the capture does not move the INPUTS, so the depth-2 race is untouched, and the fix is `StepDevInputs` as a SEAM capability, which is W4. **G1 is NOT met by W2** and is recorded owed rather than implied. **W3 DONE 2026-08-19 ([#1291](https://github.com/mudler/vllm.cpp/issues/1291)): the three remaining PLAIN BATCHED drivers migrate — `Qwen3MoeDecodeGraph`, `VoxtralDecodeGraph`, `DeepseekV2DecodeGraph` — one commit each, each with its own RED-first G2 gate.** Four of the nine drivers are now on the seam, and the six batched-driver `VLLM_CPP_CUDAGRAPH` reads `## Our baseline` item 1 counted are down to TWO, both in `qwen3_5.cpp` (W4). Each gate asserts the SEAM's counters and not the backend log, because a driver that kept its raw pair produces identical logits, an identical backend log and an identical `replay_count()`; red-first on four assertions each (`test_qwen3_moe_decode_graph_seam` 222/226, `test_voxtral_decode_graph_seam` 224/228, `test_deepseek_v2_decode_graph_seam` 224/228, all exit 1), green 3/3 each after. The G2 mutation — restoring each pre-W3 driver file, 25/102, 23/92 and 25/94 lines, each compiled clean — reddens ONLY its own gate and leaves `test_breakable_graph` 216/216 and W2's `test_qwen3_decode_graph_seam` 231/231 green. The gate harness is now SHARED (`tests/vllm/models/decode_graph_seam_harness.h`); three more copies inside `tests/` would have reproduced the duplication this row removes from `src/`. **G1 IS DELIVERED and is no longer owed** — the item W1 and W2 both carried. `tests/vllm/models/test_decode_graph_seam_g1_cuda.cpp` runs each driver COLD, CAPTURE and THREE consecutive replays against its own eager arm (selected by `max_num_reqs == 0`, so both arms are one binary on one device rather than two builds, each with its OWN device KV cache) on `thor:gpu0` through an `rc` lease — NVIDIA Thor sm_110, driver 595.78, nvcc 13.0.88, source `c905bb536`, 32 `.cu.o` objects, binary resolving `libcudart.so.13`/`libcublasLt.so.13`: **3 cases, 1600 assertions, exit 0, `5 steps x 100 logits, 0 differing, 4 replays` per driver.** The COUNT carries that claim, not the status line: with no CUDA backend the same file prints `SUCCESS!` over `assertions: 0`. Bounded honestly — synthetic tiny models rather than a checkpoint, and W2's driver shares the seam by argument rather than by measurement. **W3 also found a gate that could not fail.** The three gates' `breaks_registered == 0` mode guard is a TAUTOLOGY for any model with no registered break point, and the one production `vt::GraphBreak` in the tree is W1's in `qwen3.cpp`: flipping `kFull` to `kPiecewise` in `qwen3_moe.cpp`, one token, compiled clean and left that gate GREEN at 226/226. The mode was UNOBSERVABLE from outside a driver, so `vt::GraphBreakStats` gains `full_scopes`/`piecewise_scopes`, counted in `GraphCaptureScope`'s constructor on the ACTIVE path only, with an inert-scope control; the same flip now reds all three gates on exactly those two assertions. **NO break point is registered in these three models, deliberately**: under `kFull` it would be pass-through machinery no gate can exercise, and the break-point set is what the PIECEWISE arm needs (W4/W6). **The async decline, per driver:** Voxtral needs none (its only construction site is `VoxtralGenerateGreedy`, unreachable from the runner); Qwen3-Coder and DeepSeek carry a NEW FINDING instead — `qwen3_moe_registry.cpp:107`, `deepseek_v2_registry.cpp:106` and `glm4_moe_lite_registry.cpp:125` route an async step into a host-vector replay with no `device_token_ids` check at all, filed [#1305](https://github.com/mudler/vllm.cpp/issues/1305) with W4 as owner rather than mitigated on a measurement W3 cannot make. G5's ROCm/Tenstorrent arm is NOT discharged and moves to W5: the fleet carries no such device, so it is blocked on hardware rather than unattempted. **W4 DONE 2026-08-19 ([#1307](https://github.com/mudler/vllm.cpp/issues/1307)): the persistent device input path becomes a SEAM CAPABILITY, and the two Qwen3.5 drivers migrate.** `vt::PersistentStepInput` (`include/vt/persistent_step_input.h`, `src/vt/persistent_step_input.cpp`) binds a capture-stable device destination the DRIVER owns together with its pinned host staging block, and refreshes it in place from a host source or a DEVICE one; it owns the address-stability rule as a REFUSAL, the staging block, and the refreshing ARM as an observable (`last_source()`, `vt::StepInputStats`), and deliberately NOT the device allocation, because `Qwen3_5DecodeGraph` draws its retained inputs from a DEDICATED `DevicePool` so they never pop a block the captured forward's scratch then needs (D3). RED-first against a stub with the declared API and no guarantees: `tests/vt/test_persistent_step_input.cpp` 9 cases / 0 passed / 59 assertions / 32 failed / exit 1, GREEN after at 9/9 and 59/59; three mutations (delete the capacity refusal, make a null device source a silent no-op, collapse the host arm out of staging) each compiled clean and each reds exactly one case. `Qwen3_5DecodeGraph` and `Qwen3_5DenseDecodeGraph` open a `vt::GraphCaptureScope` over a per-slot `vt::BreakableGraph` in `kFull` and replay through it, and their `PinnedStepInputs`/`StageStepInputs` staging now runs THROUGH the capability, which is what makes it reachable rather than a class with a unit test. **Six of the nine drivers are on the seam** and `grep -rn 'std::getenv("VLLM_CPP_CUDAGRAPH")' src/` returns exactly ONE line, `src/vt/breakable_graph.cpp:61` — one switch, at last. Gate `tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp` RED-first on the MoE driver's five seam assertions (3 cases / 62 assertions / 5 failed / exit 1) and GREEN after at 7/7 and 129, G4 reading `40 values, 0 differing` per driver; G2 mutations: the whole pre-W4 file restored reds BOTH drivers (296 lines, 10 assertions), the MoE replay bypassing the container reds ONLY the MoE case (7 lines), the MoE `kFull`->`kPiecewise` flip reds ONLY its mode counters (3 lines), and deleting the `StageStepInputs` call site reds ONLY the reachability case while `test_persistent_step_input` stays 59/59 green — the difference between a class that works and a capability something reaches. **W4 FALSIFIED THIS ROW'S OWN PREMISE, which is its most important result.** This record and the spec both said the fix `qwen3.cpp`'s `DenseDecodeGraphForward`'s decline names already existed as `StepDevInputs`. It does not: `StepDevInputs` has NO token-id member, and its pinned sibling `PinnedStepInputs::token_ids` was allocated at capture, filled every step, zeroed by the poison hook, and NEVER uploaded or read — the embed runs OUTSIDE the captured region from the HOST vector in every batched driver, so **the decode graph carries no token ids to the device in ANY driver**. The dead block is removed. Consequently the DECLINE STANDS and [#1305](https://github.com/mudler/vllm.cpp/issues/1305) STAYS OPEN: W4 also read the decline's recorded cause against the tree at its own parent and found it falsified (the `DeviceTokenIdsScope` WAS live on the graph path, consumed by `EmbedInto` on all three arms at `qwen3.cpp:610,621,644 @ 338cbbfd1^`), so the measured failure is real and its mechanism is unidentified — not a state from which a refactor may retire a mitigation. The async battery was NOT run and W4 says so plainly: it needs `dgx` WITH the Qwen3-0.6B/4B checkpoints, `dgx:gpu0` was held by another session for W4's whole window, and W4's lease was `thor:gpu0`. Still NO throughput claim. W5 DONE 2026-08-19 ([#1335](https://github.com/mudler/vllm.cpp/issues/1335)): the THREE SINGLE-SHAPE drivers migrate — the DFlash draft graph, the DeepSeek V4 decode graph and the Laguna decode graph, whose own note at `laguna.cpp:2116-2119` asked for this seam by name and named V4's as the sibling that moves with it. **NINE OF NINE DRIVERS ARE ON THE SEAM and the migration is COMPLETE**: a call-shaped grep over `src/vllm/` for `BeginCapture`, `EndCaptureGraph`, `ReplayGraph` and `DestroyGraph`, with comment lines excluded, returns NOTHING. The three per-model rollback switches stay (each an A/B lever for one driver); `VLLM_CPP_CUDAGRAPH` reaches all three for the first time. **D10, the auxiliary-stream fork/join, is DISCHARGED and REACHED** — `GraphCaptureScope` owns the outstanding-fork set and joins it before `EndCaptureGraph` (port of `breakable_cuda_graph.py:353-361` plus the `wait_stream` hook `:101-153`), registered by `vt::GraphNoteFork`/`GraphNoteJoin` from `laguna.cpp:2572-2576,2612`, the only fork inside a captured region by construction. Every prior stage opened `kFull`, which has ONE segment and so no between-segments window, so the rule could not be exercised before W5 and untested machinery was not landed for it. Gated as a COUNTER and an ORDER out of one backend trace, five arms including the control where the model joins first, and two mutations (deleting the join reds only the new case on 5 assertions; making it over-fire reds it on 8). DFlash is the ONE single-shape driver gateable without a GPU, because its admission predicate names neither a device type nor a kernel registry: `test_qwen3_dflash_decode_graph_seam.cpp` RED-first 3 cases/0 passed/16 assertions/7 failed exit 1, GREEN after 3/18, and the G2 mutation reds ONLY that file while seven other suites — the driver's own `test_dflash_propose` included — stay green. **G1 RE-RUN at W5's head on `thor:gpu0`** (sm_110, driver 595.78, nvcc 13.0.88, 32 `.cu.o`, source `79dc6b5bd`) because D10 put a join on the path of EVERY segment close, so the seam changed underneath the five measured drivers: `test_decode_graph_seam_g1_cuda` 5 cases / 2066 assertions / 0 failed, each reading `0 differing, 4 replays`, plus `test_breakable_graph` 265 on the same device. **And the one thing a green build could NOT have told us was measured separately**: Laguna's capture class sits behind `#ifdef VT_MARLIN_NVFP4`, so a passing build is the SAME OBSERVATION as one that compiled the region out. `-DVT_MARLIN_NVFP4=1` is on `laguna.cpp`'s own compile command, and an undeclared identifier injected immediately after its `GraphCaptureScope` line FAILED the object build under `-Werror` (`laguna.cpp:2735`) against an rc-0 baseline, restoring to an empty diff; the identical mutation on V4 failed at `deepseek_v4.cpp:1921`. Both migrated regions are COMPILED, which retires the could-not-even-be-built half. **G1 for all three and G2 for V4 and Laguna are OWED on hardware**, per driver and per reason: V4's `CanRunResidentDecode` refuses `kCPU` and needs the four CUDA-registered kernel families, Laguna's capture class exists only under `VT_MARLIN_NVFP4`. G5's ROCm/Tenstorrent arm stays BLOCKED — the fleet is all NVIDIA — and its owner moves from W5 to the ROW. Still NO throughput claim; analysis [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) W6 DONE 2026-08-19 ([#1374](https://github.com/mudler/vllm.cpp/issues/1374)): the eligibility predicate, #1020, and the negative result on the piecewise arm. | `ACTIVE` | `CLAIM-ENG-CUDAGRAPH-BREAK-W6`; [#1163](https://github.com/mudler/vllm.cpp/issues/1163), [#1192](https://github.com/mudler/vllm.cpp/issues/1192), [#1261](https://github.com/mudler/vllm.cpp/issues/1261), [#1291](https://github.com/mudler/vllm.cpp/issues/1291), [#1307](https://github.com/mudler/vllm.cpp/issues/1307), [#1305](https://github.com/mudler/vllm.cpp/issues/1305), [#1020](https://github.com/mudler/vllm.cpp/issues/1020), [#1335](https://github.com/mudler/vllm.cpp/issues/1335), [#1374](https://github.com/mudler/vllm.cpp/issues/1374), [#1380](https://github.com/mudler/vllm.cpp/issues/1380), [#1390](https://github.com/mudler/vllm.cpp/issues/1390) | | `ENG-CUDAGRAPH-DIFFUSION` | Capture the LTX-2.5 denoise loop (fixed shapes, many identical iterations — the ideal graph target). **BLOCKED, and the blocker is ours:** the render does almost no device compute to capture | T2 | SGLang enabled BCG on this shape AFTER our pin — LTX-2 H200 two-stage 10.75s->6.90s (`d4be483efb`), SANA 1024px -26% (`6c7498113f`), SANA denoise 0.73->0.457s (`56ef810cad`). Dated events, NOT pinned evidence; their win is mostly PyTorch host tax we do not pay | NO capture at all: `grep` for capture across `src/vllm/model_executor/models/ltx2*.cpp` returns nothing | blocked by [#1024](https://github.com/mudler/vllm.cpp/issues/1024) (GPU util **exactly 0 in 321 of 347 samples**, 1.00 core of 20 held for 17+ min after staging), [#1007](https://github.com/mudler/vllm.cpp/issues/1007) (VAE decode has no device arm), [#1087](https://github.com/mudler/vllm.cpp/issues/1087) (**57-66% of wall** is ONE resolution-CONSTANT serial host phase), [#1010](https://github.com/mudler/vllm.cpp/issues/1010) (no phase-boundary log). Decision point is a MEASUREMENT of GPU-busy vs wall once device-resident, not an implementation. **The unblock order now has an owning row:** `LTX25-DEVICE-RESIDENCY` ([#1264](https://github.com/mudler/vllm.cpp/issues/1264), [ltx25-device-residency.md](specs/ltx25-device-residency.md)) stages those defects W0-W6 and carries this decision point as its W7 — if the loop comes back GPU-bound, #1164 closes as a refutation the way [#1161](https://github.com/mudler/vllm.cpp/issues/1161) closed prefill capture | [sglang-breakable-cuda-graph.md](specs/sglang-breakable-cuda-graph.md) | `INVENTORIED` | [#1164](https://github.com/mudler/vllm.cpp/issues/1164) | | `ENG-BATCH-INVARIANT` | Opt-in deterministic execution across scheduler batch sizes (`VLLM_BATCH_INVARIANT=1`): batch-invariant matmul/norm/attention/collectives plus persistent-scheduler NVFP4; production default remains off | T1 | default/env `vllm/envs.py:89,576-578`; initialization `vllm/v1/worker/gpu_worker.py:1262`; NVFP4 dispatch `csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu:212-220`; suite fixture `tests/v1/determinism/conftest.py:9-12`; operator/e2e `tests/v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py`, `tests/v1/determinism/test_nvfp4_batch_invariant.py` @ `702f481` | - | [W3-C3R executed contract](specs/nvfp4-persistent-plan-cache.md#w3-c3r-batch-shape-localization-and-gate-correction-2026-07-13): production-default ours and vLLM both change outputs across batch shapes; no local opt-in implementation is claimed | `planned: specs/batch-invariant-execution.md` | `INVENTORIED` | - | -| `ENG-ASYNC-SCHED` | Async/overlap scheduling (AsyncScheduler placeholders + depth-2 batch-queue step + async D2H on a copy stream); vLLM's DEFAULT at the pin — mirror obligation per B3. **Host-side machinery + runner device-input half + sampler-OUTPUT half LANDED + CPU-gated (2026-07-16):** `AsyncScheduler` placeholder accounting, `step_with_batch_queue` depth-2, `ResolveAsyncScheduling` default-ON-when-compatible + `MaxConcurrentBatches`, `VT_ASYNC_SCHED` rollback; the runner device-input path `combine_sampled_and_draft_tokens`; PLUS the sampler-OUTPUT half — `vt::Backend` event/pinned primitives (`AllocPinned`/events, CUDA cudaHostAlloc+cudaEvent, CPU sync-degeneration), `AsyncGPUModelRunnerOutput` (device sampled-id snapshot → non-blocking D2H on a copy queue + event; `get_output()` waits only that event; MAIN queue never blocked), `Sampler::forward(sampled_ids_out)` device-resident greedy, `GPUModelRunner::sample_tokens_async` + `runner_supports_async`, and the `Executor`+`step_with_batch_queue` seam resolving `get_output()` at CONSUME time. All behind `VT_ASYNC_RUNNER`/`set_async_input_combine`, default OFF. Sync path byte-identical (placeholder sites INERT while count 0; combine off; `sample_tokens_async` degenerates to sync when async off; `sampled_ids_out=nullptr`). **ENABLE-FLIP LANDED + CPU-gated (2026-07-16):** (1) `LoadedEngine` now reorders `runner_` before the scheduler and builds an `AsyncScheduler` + `max_concurrent_batches=2` when `ResolveAsyncScheduling(runner_.runner_supports_async())` resolves ON (else the byte-identical synchronous `Scheduler` + depth-1); the resolved mcb threads into `AsyncLLM`→`EngineCoreProc` (`step_with_batch_queue`) and the "Asynchronous scheduling is enabled/disabled" log mirrors vLLM for A/B audit; (2) the device combine/scatter kernel (`_combine_sampled_and_draft_tokens_kernel` + last_sampled scatter) is ported to CUDA (`src/vt/cuda/cuda_combine_tokens.cu`), main-stream-ordered on the CUDA async path so it DELETES `sample_tokens_async`'s pre-scatter `Synchronize`; the CPU backend keeps the host loop. `VT_ASYNC_RUNNER=1` engages full W3; `VT_ASYNC_SCHED=0` is the same-binary rollback. Production default (no env) stays synchronous byte-identical. **FULL W3 DGX proof RAN twice** — `f086b64` (5/5 gates PASS; c16 TPOT −5.4 ms WIN, tput neutral, TTFT +36 % = Little's-law repayment) and the 2026-07-16 re-proof on the THROUGHPUT-lever fix (persistent pooled sampled-id/pinned buffers + `Sampler` greedy scratch removing ALL per-step `cudaMalloc`/`cudaFree`/`cudaHostAlloc`/event-create from the sampled-id path, incl. the overlap-killing `cudaFree` inside `get_output`; mirrors `gpu_model_runner.py:873-878` + `async_utils.py:12-70`): token-exactness **6/6 PASS**, interleaved c16 **tput −0.32 % (gate ≥+1.5 % FAILS), TPOT −4.95 ms retained, TTFT +34.8 %** — the allocator lever is REFUTED as the tput unlock (≤0.1 % of a ~165 ms c16 step). **DEFAULT FLIPPED ON 2026-07-17** (`VT_ASYNC_RUNNER` default ON via the pure `AsyncRunnerFlagIsOn` predicate, mirroring `vllm/config/vllm.py:992-1044`): the discriminator (`6ea7856`) proved vLLM's own async pays the identical +26–31 % TTFT / −0.7 to −0.9 % tput / −2.6 to −4.3 ms TPOT envelope and W3-ON nets positive (both binding ITL-tail anomalies flip to PASS), so the "needs a throughput lever" ship-gate is RETIRED — W3 is a parity/mirror obligation with a tails+TPOT win. The flip is TOKEN-NEUTRAL (async-ON ≡ async-OFF bit-identical on DGX). `VT_ASYNC_RUNNER=0` = runner-level rollback, `VT_ASYNC_SCHED=0` = scheduler-level rollback. TTFT means rise into vLLM's async envelope BY DESIGN — the next binding grid runs async by default and its TTFT must NOT be misread as a regression. **ROBUSTNESS FIX 2026-07-20 (`discard_request_mask`):** the runner was missing vLLM's `discard_request_mask`, so `GPUModelRunner` emitted a sampled token for prefill-CHUNK requests too; under async this drained a `num_output_placeholders` never reserved (the `is_prefill_chunk` path adds none) → the `async_scheduler.cpp` `num_output_placeholders >= 0` assertion aborted on c8 + short-output (chunked prefill + preemption). FIX mirrors vLLM: `execute_model` computes `exec_state_.discard[i] = seq_len < num_tokens` (`gpu_model_runner.py:2048`); `sample_tokens` clears those rows to empty (`outputs.py:303`), the async path passes `invalid_req_indices` to `AsyncGPUModelRunnerOutput::get_output` (`gpu_model_runner.py:3625` + `outputs.py:303`). Scheduler UNCHANGED (assertion kept — it was correct once the runner honors `scheduler.py:1888-1890`). Sync/non-chunked decode byte-identical (mask all-zero); DGX 27B 235/235 + 35B 315/315, `vllm-bench` c8+short-output+chunked+kv-pressure no longer crashes, memcheck 0. Ledger [parity-ledger.md](parity-ledger.md) 2026-07-20 row | T1 | `vllm/v1/core/sched/async_scheduler.py:12`; `vllm/config/vllm.py:490,990,1038`; `vllm/v1/engine/core.py:519`; `vllm/v1/worker/gpu/input_batch.py:304-406`; `vllm/v1/worker/gpu/async_utils.py:12-70`; `vllm/v1/worker/gpu/gpu_model_runner.py:242-332`; `vllm/v1/outputs.py:298-307` | `src/vllm/v1/core/sched/async_scheduler.cpp:10,45`; placeholder plumbing `src/vllm/v1/core/sched/scheduler.cpp:148,164,605`; `src/vllm/v1/engine/core.cpp:91` (`step_with_batch_queue`, async-output seam); `src/vllm/v1/engine/core_proc.cpp:32,46`; config `include/vllm/config/scheduler.h:117,165,188`, `src/vllm/config/scheduler.cpp:12`; `include/vllm/v1/request.h:187`; runner input leaf `src/vllm/v1/worker/gpu/prepare_inputs.cpp`, `src/vllm/v1/worker/gpu/input_batch.cpp`; runner output leaf `include/vt/backend.h`+`src/vt/backend.cpp`+`src/vt/cuda/cuda_backend.cu` (event/pinned), `include/vllm/v1/worker/gpu/async_output.{h,cpp}` (`AsyncGPUModelRunnerOutput`), `src/vllm/v1/sample/sampler.cpp` (`sampled_ids_out`), `src/vllm/v1/worker/gpu/runner.cpp` (`sample_tokens_async`/`runner_supports_async`), `src/vllm/v1/executor/executor.cpp`+`include/vllm/v1/worker/gpu/model_runner_base.h` (async seam); enable-flip `include/vllm/entrypoints/model_loader.h`+`src/vllm/entrypoints/model_loader.cpp` (`runner_` before scheduler, `ResolveAsyncEnabled`/`MakeScheduler`, `AsyncScheduler`+mcb=2, log), `include/vllm/v1/engine/async_llm.h`+`src/vllm/v1/engine/async_llm.cpp` (mcb param → `EngineCoreProc`); device kernel `include/vt/cuda/combine_tokens.h`+`src/vt/cuda/cuda_combine_tokens.cu`, wired `src/vllm/v1/worker/gpu/runner.cpp` (CUDA combine/scatter branch removes the pre-sync) | `tests/vllm/v1/test_async_scheduler.cpp:1` (6 cases, 54 asserts; RED vs base Scheduler 2/6 fail); depth-2 engine cycle `tests/vllm/v1/test_engine_core_proc.cpp:479` (mcb=2, async-output seam); config resolution `tests/vllm/test_scheduler_config.cpp:75`; enable-flip construction matrix `tests/vllm/entrypoints/test_loaded_engine_dense.cpp` (runner×VT_ASYNC_SCHED → scheduler type + mcb; RED = un-flipped engine, 3/3 ON-arm asserts fail); runner input leaf `test_combine_tokens.cpp` (RED = stale → 5/7 fail), `test_input_batch.cpp`, `test_runner.cpp` (async-ON≡sync); output leaf `tests/vt/test_backend.cpp` (event/pinned contract), `tests/vllm/v1/worker/test_async_output.cpp` (materialize/flush/snapshot; RED = +1 splice), `test_runner.cpp` (`sample_tokens_async` decode ≡ sync); full CPU ctest 111/111, tools 164/164. Prior diagnostic `3812d8` six-leg control: total **1.002153×**, TTFT **0.862159×**, no GPU-time reduction (neutral for speed). **DEFAULT-FLIP (2026-07-17):** new pure CPU flag test [test_async_runner_flag.cpp](../tests/vllm/v1/worker/test_async_runner_flag.cpp) (11 asserts, default-ON/'0'-off); construction matrix [test_loaded_engine_dense.cpp](../tests/vllm/entrypoints/test_loaded_engine_dense.cpp) INVERTED (default → AsyncScheduler+mcb=2; RED verified 5 asserts fail vs un-flipped). CPU clean `-Werror` rebuild, full serial ctest **116/116**, tools **164/164**. **DGX re-confirmation** (evidence `dgx:~/work/vllm.cpp-async-flip`, CUTLASS+FA2 hard-verified, one flock): shipping default (async ON + RMSNorm-fast OFF) → **27B 235/235 + 35B 315/315** with the "Asynchronous scheduling is enabled (mcb=2)" log, and both rollback arms (`VT_ASYNC_RUNNER=0`, `VT_ASYNC_SCHED=0`) 235/235 + 315/315 log "disabled"; async arms BIT-IDENTICAL (token-neutral). Closing record [parity-ledger.md#L502](parity-ledger.md#L502) | [async-serving.md](specs/async-serving.md) | `DONE` | `6ea7856` | +| `ENG-ASYNC-SCHED` | Async/overlap scheduling (AsyncScheduler placeholders + depth-2 batch-queue step + async D2H on a copy stream); vLLM's DEFAULT at the pin — mirror obligation per B3. **Host-side machinery + runner device-input half + sampler-OUTPUT half LANDED + CPU-gated (2026-07-16):** `AsyncScheduler` placeholder accounting, `step_with_batch_queue` depth-2, `ResolveAsyncScheduling` default-ON-when-compatible + `MaxConcurrentBatches`, `VT_ASYNC_SCHED` rollback; the runner device-input path `combine_sampled_and_draft_tokens`; PLUS the sampler-OUTPUT half — `vt::Backend` event/pinned primitives (`AllocPinned`/events, CUDA cudaHostAlloc+cudaEvent, CPU sync-degeneration), `AsyncGPUModelRunnerOutput` (device sampled-id snapshot → non-blocking D2H on a copy queue + event; `get_output()` waits only that event; MAIN queue never blocked), `Sampler::forward(sampled_ids_out)` device-resident greedy, `GPUModelRunner::sample_tokens_async` + `runner_supports_async`, and the `Executor`+`step_with_batch_queue` seam resolving `get_output()` at CONSUME time. All behind `VT_ASYNC_RUNNER`/`set_async_input_combine`, default OFF. Sync path byte-identical (placeholder sites INERT while count 0; combine off; `sample_tokens_async` degenerates to sync when async off; `sampled_ids_out=nullptr`). **ENABLE-FLIP LANDED + CPU-gated (2026-07-16):** (1) `LoadedEngine` now reorders `runner_` before the scheduler and builds an `AsyncScheduler` + `max_concurrent_batches=2` when `ResolveAsyncScheduling(runner_.runner_supports_async())` resolves ON (else the byte-identical synchronous `Scheduler` + depth-1); the resolved mcb threads into `AsyncLLM`→`EngineCoreProc` (`step_with_batch_queue`) and the "Asynchronous scheduling is enabled/disabled" log mirrors vLLM for A/B audit; (2) the device combine/scatter kernel (`_combine_sampled_and_draft_tokens_kernel` + last_sampled scatter) is ported to CUDA (`src/vt/cuda/cuda_combine_tokens.cu`), main-stream-ordered on the CUDA async path so it DELETES `sample_tokens_async`'s pre-scatter `Synchronize`; the CPU backend keeps the host loop. `VT_ASYNC_RUNNER=1` engages full W3; `VT_ASYNC_SCHED=0` is the same-binary rollback. Production default (no env) stays synchronous byte-identical. **FULL W3 DGX proof RAN twice** — `f086b64` (5/5 gates PASS; c16 TPOT −5.4 ms WIN, tput neutral, TTFT +36 % = Little's-law repayment) and the 2026-07-16 re-proof on the THROUGHPUT-lever fix (persistent pooled sampled-id/pinned buffers + `Sampler` greedy scratch removing ALL per-step `cudaMalloc`/`cudaFree`/`cudaHostAlloc`/event-create from the sampled-id path, incl. the overlap-killing `cudaFree` inside `get_output`; mirrors `gpu_model_runner.py:873-878` + `async_utils.py:12-70`): token-exactness **6/6 PASS**, interleaved c16 **tput −0.32 % (gate ≥+1.5 % FAILS), TPOT −4.95 ms retained, TTFT +34.8 %** — the allocator lever is REFUTED as the tput unlock (≤0.1 % of a ~165 ms c16 step). **DEFAULT FLIPPED ON 2026-07-17** (`VT_ASYNC_RUNNER` default ON via the pure `AsyncRunnerFlagIsOn` predicate, mirroring `vllm/config/vllm.py:992-1044`): the discriminator (`6ea7856`) proved vLLM's own async pays the identical +26–31 % TTFT / −0.7 to −0.9 % tput / −2.6 to −4.3 ms TPOT envelope and W3-ON nets positive (both binding ITL-tail anomalies flip to PASS), so the "needs a throughput lever" ship-gate is RETIRED — W3 is a parity/mirror obligation with a tails+TPOT win. The flip is TOKEN-NEUTRAL (async-ON ≡ async-OFF bit-identical on DGX). `VT_ASYNC_RUNNER=0` = runner-level rollback, `VT_ASYNC_SCHED=0` = scheduler-level rollback. TTFT means rise into vLLM's async envelope BY DESIGN — the next binding grid runs async by default and its TTFT must NOT be misread as a regression. **ROBUSTNESS FIX 2026-07-20 (`discard_request_mask`):** the runner was missing vLLM's `discard_request_mask`, so `GPUModelRunner` emitted a sampled token for prefill-CHUNK requests too; under async this drained a `num_output_placeholders` never reserved (the `is_prefill_chunk` path adds none) → the `async_scheduler.cpp` `num_output_placeholders >= 0` assertion aborted on c8 + short-output (chunked prefill + preemption). FIX mirrors vLLM: `execute_model` computes `exec_state_.discard[i] = seq_len < num_tokens` (`gpu_model_runner.py:2048`); `sample_tokens` clears those rows to empty (`outputs.py:303`), the async path passes `invalid_req_indices` to `AsyncGPUModelRunnerOutput::get_output` (`gpu_model_runner.py:3625` + `outputs.py:303`). Scheduler UNCHANGED (assertion kept — it was correct once the runner honors `scheduler.py:1888-1890`). Sync/non-chunked decode byte-identical (mask all-zero); DGX 27B 235/235 + 35B 315/315, `vllm-bench` c8+short-output+chunked+kv-pressure no longer crashes, memcheck 0. Ledger [parity-ledger.md](parity-ledger.md) 2026-07-20 row | T1 | `vllm/v1/core/sched/async_scheduler.py:12`; `vllm/config/vllm.py:490,990,1038`; `vllm/v1/engine/core.py:519`; `vllm/v1/worker/gpu/input_batch.py:304-406`; `vllm/v1/worker/gpu/async_utils.py:12-70`; `vllm/v1/worker/gpu/gpu_model_runner.py:242-332`; `vllm/v1/outputs.py:298-307` | `src/vllm/v1/core/sched/async_scheduler.cpp:10,45`; placeholder plumbing `src/vllm/v1/core/sched/scheduler.cpp:148,164,605`; `src/vllm/v1/engine/core.cpp:137` (`step_with_batch_queue`, async-output seam); `src/vllm/v1/engine/core_proc.cpp:32,46`; config `include/vllm/config/scheduler.h:117,165,188`, `src/vllm/config/scheduler.cpp:12`; `include/vllm/v1/request.h:187`; runner input leaf `src/vllm/v1/worker/gpu/prepare_inputs.cpp`, `src/vllm/v1/worker/gpu/input_batch.cpp`; runner output leaf `include/vt/backend.h`+`src/vt/backend.cpp`+`src/vt/cuda/cuda_backend.cu` (event/pinned), `include/vllm/v1/worker/gpu/async_output.{h,cpp}` (`AsyncGPUModelRunnerOutput`), `src/vllm/v1/sample/sampler.cpp` (`sampled_ids_out`), `src/vllm/v1/worker/gpu/runner.cpp` (`sample_tokens_async`/`runner_supports_async`), `src/vllm/v1/executor/executor.cpp`+`include/vllm/v1/worker/gpu/model_runner_base.h` (async seam); enable-flip `include/vllm/entrypoints/model_loader.h`+`src/vllm/entrypoints/model_loader.cpp` (`runner_` before scheduler, `ResolveAsyncEnabled`/`MakeScheduler`, `AsyncScheduler`+mcb=2, log), `include/vllm/v1/engine/async_llm.h`+`src/vllm/v1/engine/async_llm.cpp` (mcb param → `EngineCoreProc`); device kernel `include/vt/cuda/combine_tokens.h`+`src/vt/cuda/cuda_combine_tokens.cu`, wired `src/vllm/v1/worker/gpu/runner.cpp` (CUDA combine/scatter branch removes the pre-sync) | `tests/vllm/v1/test_async_scheduler.cpp:1` (6 cases, 54 asserts; RED vs base Scheduler 2/6 fail); depth-2 engine cycle `tests/vllm/v1/test_engine_core_proc.cpp:479` (mcb=2, async-output seam); config resolution `tests/vllm/test_scheduler_config.cpp:75`; enable-flip construction matrix `tests/vllm/entrypoints/test_loaded_engine_dense.cpp` (runner×VT_ASYNC_SCHED → scheduler type + mcb; RED = un-flipped engine, 3/3 ON-arm asserts fail); runner input leaf `test_combine_tokens.cpp` (RED = stale → 5/7 fail), `test_input_batch.cpp`, `test_runner.cpp` (async-ON≡sync); output leaf `tests/vt/test_backend.cpp` (event/pinned contract), `tests/vllm/v1/worker/test_async_output.cpp` (materialize/flush/snapshot; RED = +1 splice), `test_runner.cpp` (`sample_tokens_async` decode ≡ sync); full CPU ctest 111/111, tools 164/164. Prior diagnostic `3812d8` six-leg control: total **1.002153×**, TTFT **0.862159×**, no GPU-time reduction (neutral for speed). **DEFAULT-FLIP (2026-07-17):** new pure CPU flag test [test_async_runner_flag.cpp](../tests/vllm/v1/worker/test_async_runner_flag.cpp) (11 asserts, default-ON/'0'-off); construction matrix [test_loaded_engine_dense.cpp](../tests/vllm/entrypoints/test_loaded_engine_dense.cpp) INVERTED (default → AsyncScheduler+mcb=2; RED verified 5 asserts fail vs un-flipped). CPU clean `-Werror` rebuild, full serial ctest **116/116**, tools **164/164**. **DGX re-confirmation** (evidence `dgx:~/work/vllm.cpp-async-flip`, CUTLASS+FA2 hard-verified, one flock): shipping default (async ON + RMSNorm-fast OFF) → **27B 235/235 + 35B 315/315** with the "Asynchronous scheduling is enabled (mcb=2)" log, and both rollback arms (`VT_ASYNC_RUNNER=0`, `VT_ASYNC_SCHED=0`) 235/235 + 315/315 log "disabled"; async arms BIT-IDENTICAL (token-neutral). Closing record [parity-ledger.md#L502](parity-ledger.md#L502) | [async-serving.md](specs/async-serving.md) | `DONE` | `6ea7856` | | `ENG-PRIORITY-SCHED` | Priority request queue + policy + priority preemption + `priority` plumbing (Request/EngineCoreRequest/OpenAI field); W4 of the async-serving block. Default stays FCFS. GATING: full CPU tier green (93/93; 12 ported priority-scheduler cases + 14 priority-queue cases incl. the seeded random property test); ~~GPU G1 (both greedy engine gates, priority-vs-fcfs token-exactness) deferred to the next GPU-idle window — GPU held by the `SERVE-GATE-ONLINE` campaign~~. **BLOCKER CORRECTED 2026-08-12 ([#534](https://github.com/mudler/vllm.cpp/issues/534)): the GPU is not what is stopping this, and G1 cannot be "rerun" because it does not exist.** `--scheduling-policy priority` is plumbed to the production server (`src/vllm/entrypoints/openai/server_main.cpp:408-411,672-673` -> `SchedulerPolicyFromString` -> `SchedulerConfig::policy`), but NO priority-vs-fcfs gate exists at the ENGINE/MODEL level. What exists is the scheduler-unit tier (`test_scheduler.cpp:674,916` admission/preemption ordering, `test_request_queue.cpp` heap ordering) plus C-ABI wire-name validation (`tests/capi/test_capi.cpp:1159`); nothing anywhere drives a real engine with `policy=kPriority` and compares token streams against the fcfs arm. A next owner writes that gate RED-first, then runs it. The struck deferral is a 2026-07-10 scheduling note that five successive `SERVE-GATE-ONLINE` bindings (`3f256ab`, `246a23c`, `a875397`, `f0fb727`, `9ecd9d0`) have since expired. This is the ONLY genuinely open leaf of the `ROAD-V1-C6` async-serving block | T1 | `vllm/v1/core/sched/request_queue.py:131,201`; `vllm/v1/core/sched/scheduler.py:546`; `vllm/config/scheduler.py:109`; `tests/v1/core/test_scheduler.py:2382,2978`; `tests/v1/core/test_priority_scheduler_random.py:1` | `src/vllm/v1/core/sched/request_queue.cpp:101,186`; `src/vllm/v1/core/sched/scheduler.cpp:178`; `src/vllm/v1/request.cpp:92`; `src/vllm/config/scheduler.cpp:21` | `tests/vllm/v1/test_scheduler.cpp:674,916`; `tests/vllm/v1/test_request_queue.cpp:238,429` | [async-serving.md](specs/async-serving.md) | `GATING` | - | | `ENG-PARTIAL-PREFILL` | Concurrent partial-prefill and long-prompt limits | T1 | `vllm/config/scheduler.py:70-80` | - | - | `planned: specs/partial-prefill-concurrency.md` | `INVENTORIED` | - | | `ENG-BATCH-QUEUE` | Pipelined `step_with_batch_queue` | T1 | `vllm/v1/engine/core.py:519` | - | - | `planned: specs/batch-queue-step.md` | `INVENTORIED` | - | @@ -74,7 +74,7 @@ forensics: roadmap_v1.md and the parity ledger. | `ENG-DBO-UBATCH` | DBO and ubatch overlap | T2 | `vllm/config/parallel.py:208,524` | - | - | `planned: specs/dbo-ubatch.md` | `INVENTORIED` | - | | `ENG-MOE-SHARED-AUX` | MoE shared-expert MLP on an aux CUDA stream concurrent with the routed-expert router/align/grouped-GEMMs (mirror vLLM's decode overlap; the largest remaining 35B c1/c2 engine lever). Fork the shared MLP onto a 2nd persistent per-device stream, join before the combine → byte-identical to serial (independent shared/routed paths both complete before combine; overlap changes WHEN not WHAT). Gated `T <= threshold` decode + CUDA. The aux stream draws scratch from a SEPARATE `AuxPool` so the concurrent main-stream routed allocations never share a live block with it (the `DevicePool` reuse invariant is single-stream ordering; vLLM sidesteps this with its stream-aware caching allocator's `record_stream`). `VT_MOE_SHARED_AUX_STREAM` **DEFAULT ON** (`=0` rollback); `VT_MOE_SHARED_AUX_THRESHOLD` (default 128; GB10 48-SM calibration). Captured in the decode CUDA-graph via the fork/join event edges (`ThreadLocal` capture, no abort). Only the committed Marlin MoE decode path; wmma fallback/CPU/GGUF and 27B dense unaffected | T1 | `vllm/model_executor/layers/fused_moe/runner/shared_experts.py:99-104,125-142`; `vllm/utils/multi_stream_utils.py:20-58` (`maybe_execute_in_parallel`, TRT-LLM port); `vllm/utils/torch_utils.py:736-756` (`aux_stream`); `vllm/envs.py:260` (threshold 256) | fork/join `src/vllm/model_executor/models/qwen3_5.cpp:3999,4114` (`MoeBlockFusedMarlinCuda`); aux stream+events `src/vllm/model_executor/models/qwen3_5.cpp:3575,3581` (`MoeAuxStream`/`MoeAuxStreamFor`); predicates `:3553,3560`; aux-pool isolation `:496,3538` (`AuxPool`/`ActivePool`/`ActivePoolScope`) + `DBuf pool_` routing `:645` | **DGX (prod flags, one flock):** overlap ON==OFF BYTE-IDENTICAL — `tests/parity/test_qwen36_paged_engine.cpp:116` 35B **315/315** + `tests/parity/test_qwen27_paged_engine.cpp` 27B **235/235** under `VT_MOE_SHARED_AUX_STREAM`∈{0,1}; captured-vs-eager (`VLLM_CPP_CUDAGRAPH=0`, ON) 315/315; shipping default (no env) 315/315+235/235, rollback `=0` 315/315+235/235; `compute-sanitizer memcheck` (default ON, captured) 0 errors; in-situ interleaved TPOT A/B (drop cold rep1) c1 −5.6% / c2 −2.7% / c4 −3.7% / c8 −3.4% / c16 −1.6% / c32 −1.5% (WINS every conc, zero regression); ledger [parity-ledger.md](parity-ledger.md) 2026-07-19 row | [moe-shared-aux-stream.md](specs/moe-shared-aux-stream.md) | `ANCHOR-BACKFILL` | `CLAIM-MOE-SHARED-AUX-1` | | `ENG-RUNNER-MODELSHAPE` | **Runner is model-shape-agnostic over the KV-cache group structure** — the extensibility deliverable the first additive-model bring-up (Qwen3 dense) forced. Before W1 the `GPUModelRunner` had only ever executed the Qwen3.6 HYBRID topology and hardcoded it in two places: (#1) the KV-buffer alloc loop indexed `config_.layer_types[l]`, out-of-bounds on a pure-dense model's EMPTY `layer_types`; (#2) each `execute_model` step unconditionally built the GDN metadata (`gather_block_table(gdn_group_id_)` / `remap_gdn_state_slots` / `GDNAttentionMetadataBuilder`), which reads `block_table[-1]` when there is no mamba group. W1 drives both off the resolved KV-group structure — a model-agnostic `has_mamba_group` / `gdn_group_id_ >= 0` predicate (NOT a model-name check): empty/absent `layer_types` ⇒ all full-attention; no mamba group ⇒ the whole GDN metadata/state path is skipped and `gdn_meta` stays default-empty. A full-attention-only KV config (one FA group, no MambaSpec) now allocates + steps cleanly; the hybrid gate models keep their GDN group so their path is BYTE-IDENTICAL. This is a one-time generalization: every future dense/non-hybrid arch (Llama, Mistral) now adds new-files-only, zero further runner edits. **PER-LAYER KV head_dim extension (Gemma-4 G1b, 2026-07-28, `CLAIM-GEMMA4-G1B`):** the runner's full-attn alloc/view loops now consume an OPTIONAL `KVCacheConfig::per_layer_attn_specs` (index == layer) so a HETEROGENEOUS-head_dim model (Gemma-4: sliding 256 / global 512, same num_kv_heads) sizes each non-GDN layer's paged KV + PagedKvCache view from its OWN spec. The field is EMPTY for every uniform-KV model ⇒ the loop collapses to the single group spec ⇒ byte-identical allocation/view/indexing/dispatch (same additive-identical property as the model-shape generalization above). Block table / KV manager / scheduler stay head_dim-independent (num_blocks + block_size, uniform) so no per-group block table is introduced | T0 | model-agnostic runner drives off `kv_cache_config.kv_cache_groups` — `vllm/v1/worker/gpu/model_runner.py` `initialize_kv_cache` / attention-metadata build (per-group, no hardcoded hybrid) @ `e24d1b24` | `src/vllm/v1/worker/gpu/runner.cpp:458-470` (alloc loop: `has_mamba_group && !layer_types.empty()` gate) + `:651-680` (GDN metadata build gated on `gdn_group_id_ >= 0`, default-empty `gdn_meta` otherwise); per-layer KV head_dim: `include/vllm/v1/kv_cache_interface.h` (`KVCacheConfig::per_layer_attn_specs`) consumed in `src/vllm/v1/worker/gpu/runner.cpp` `initialize_kv_cache` (per-layer `FaDims` alloc+view), published by `src/vllm/model_executor/models/gemma4_registry.cpp` (`MakeGemma4ForConditionalGenerationKVCache`); the full-attention-only KV spec that exercises the base path `src/vllm/model_executor/models/qwen3_dense.cpp` (`MakeQwen3ForCausalLMKVCache`) | `tests/vllm/v1/worker/test_runner.cpp:1129` — "full-attention-only KV config allocates without the GDN path" + "full-attention-only step skips GDN metadata build (no OOB)" (RED→GREEN: both SIGSEGV pre-generalization; GREEN post). Behaviour-preservation gate: DGX **27B 235/235 + 35B 315/315 UNCHANGED** under the fix; per-layer-KV inertness: full CPU runner/KV suite green + **OLMo-2 SACRED GPU re-gate 16/16 UNCHANGED**; heterogeneous path proven by **Gemma-4 E4B STRICT 32/32** (`tests/parity/test_gemma4_paged_engine.cpp`); ASan/UBSan clean on the affected paths | [first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md) §3 (seam gaps #1/#2), §6 (W1); [gemma4-multimodal.md](specs/gemma4-multimodal.md) §G1b | `ACTIVE` | `CLAIM-MODEL-QWEN3-DENSE` | -| `ENG-MM-INPUT-PIPELINE` | **Multimodal INPUT pipeline + encoder-cache engine seam (M1), INERT when no mm input.** The C++ mirror of `vllm/multimodal/`: `MultiModalKwargs`/`MultiModalFeatureSpec`/`MultiModalInputs`, the `MultiModalHasher` mm-hash (blake3), the Qwen3-VL image processor (smart_resize + fused rescale/normalize + patchify -> `pixel_values`+`image_grid_thw`) and placeholder-token expansion, plus the `EncoderCacheManager` (+`ComputeMmEncoderBudget`) and the LMCache `extra_keys` seam. Additive `mm_features` carried on `Request`/`EngineCoreRequest`; with NO mm input every field is empty and every path is byte-identical to the text engine. Processor output is BIT/BYTE-identical to the vLLM 0.25.0 oracle (M0 fixture). Does NOT build the vision tower / embed-merge (M2). **SERVING wiring (ROAD-V1-MM `MM-SERVE-ENGINE`, 2026-07-28, `CLAIM-MM-SERVING-W2`):** the OpenAI server now carries the parsed `MultiModalInputs` into the engine — additive `LLMEngine`/`AsyncLLM` `add_request(MultiModalInputs)`+`generate(MultiModalInputs)` overloads via `InputProcessor::process_inputs_mm` (mirror `input_processor.py:333-379`, empty mm_features == the tokens path), the chat-template placeholder-STRING helpers (`get_placeholder_str`/`_add_placeholder` mirror), and the serving_chat `MultiModalChatFn` seam (default unset ⇒ text byte-identical). **SEAM BODY (ROAD-V1-MM `MM-SERVE-E2E` W3, 2026-07-28, `CLAIM-MM-SERVING-E2E`):** `MakeQwen3VLImageChatFn` (chat_mm.cpp) is the seam body the server sets — messages → marker-inject → chat template → `EncodeWithSpecialTokens` (the single image_pad marker → one image_token_id) → `RouteImageRgb` EXPAND to 196 image tokens + mm_features; wired in `examples/server/main.cpp` (guarded on `preprocessor_config.json`; text-only unset ⇒ byte-identical). Gated `test_chat_mm` 8/8 + `test_openai_serving` (seam invoked + routed). **ENGINE MM-FORWARD LANDED (ROAD-V1-MM `MM-SERVE-E2E`, 2026-07-28, `CLAIM-ENGINE-MM-FORWARD`):** the engine model runner now HAS an mm forward — `ModelForwardInput` gains an ADDITIVE default-nullopt `std::optional mm` (merged inputs_embeds + 3-D MRoPE positions + DeepStack, borrowed handles; nullopt-for-text ⇒ shared runner path byte-identical BY CONSTRUCTION), `Qwen3VLForConditionalGeneration` is `REGISTER_VLLM_MODEL`-registered (`qwen3_vl_registry.cpp`), and the registered forward FOLDS the M2c decode into `ModelRegistry::Forward` via the SHARED `Qwen3VLForwardStepLastLogits` (`Qwen3VLGenerateGreedyViaRegistry` drives every step through the registry). GPU token-exact gate `test_qwen3vl_registry_e2e` (image→text THROUGH `ModelRegistry::Forward` == M2c golden 32/32 STRICT, dgx.casa GB10); text inertness `test_runner` 16/16 + `test_scheduler` 36/36 + `test_model_registry` 24/24 + `test_chat_mm` 8/8 + `test_openai_serving` 41/41 all green. RESIDUAL: the FULL in-runner scheduler-fed tower run (batched-loop mm building the field from staged encoder outputs) + the real server `/v1/chat/completions` GPU e2e — recipe in `specs/mm-serving.md`. **INPUT LIMITS L1 LANDED (#607, 2026-08-13):** the per-modality `limit_per_prompt` + `GetLimitPerPrompt` precedence (`language_model_only` ⇒ 0 BEFORE the map, else the map, else 999) and the refusal that gives those numbers effect — `AllowedMmLimits` folding by `min()` against the model's own ceiling, `ValidateNumItems` with upstream's exact message, and both call sites with the `enable_mm_embeds` escape. NO serve surface and NO live call site: nothing constructs a `MultiModalConfig` on a request yet, which is L2's. **INPUT LIMITS L2 LANDED (#607, #686, 2026-08-14):** the flags (`--[no-]language-model-only`, `--limit-mm-per-prompt ''`; `arg_utils.py:555-556,1276-1279,1691-1692` over `ParseLimitMmPerPromptJson`, the port of `multimodal.py:212-236` + the DummyOptions dataclasses `:17-43`), the C-ABI fields (`vllm_model_params.language_model_only`/`.limit_mm_per_prompt`, **ABI v19**), and the LIVE CALL SITE: `ValidateChatMmLimits` (`chat_utils.py:648-662`) runs as step 0 of `MakeQwen3VLImageChatFn` over a `BaseProcessingInfo` folding `LoadedEngine::mm_config()` with `Qwen3VLChatSupportedMmLimits() == {"image": 1}` — the seam's own ceiling, which is the `min()` fold operand #686 recorded as undeclared. A three-image request is now HTTP 400 with upstream's message rather than an opaque 500 / a truncated answer. NOT claimed AT L2: any memory win — L3 gates tower construction, and its saving is MEASURED only for `Qwen3-VL-4B-Instruct` (host RSS at load, not VRAM, and half 1 of a two-half gate); `muse-glimmer` and the `--mmproj` site stay unmeasured. Still unwired: the `process_inputs_mm` call site (`context.py:461`), blocked on the per-model `get_supported_mm_limits()` hook. **TOWER SKIP L3 (#607):** the predicate `SkipTowerForModalities` (the port of `interfaces.py:293` — `all(...) == 0` over the tower's OWN modality set, ALL not ANY, reached through `GetLimitPerPrompt` so `--limit-mm-per-prompt '{"image":0,"video":0}'` skips exactly as `--language-model-only` does), the borrow that carries the engine's limits into a load (`ModelSource::multimodal`), and the observable that says what was freed (`LoadedModel::skipped_towers` / `LoadedEngine::skipped_towers`, printed once by the server). All THREE production tower loads are gated: `LoadQwen3VLWeights`, `LoadMuseGlimmerForConditionalGenerationWeights`, and — added in the L3 repair wave, missed by the first cut because it lives in the entrypoint rather than in a model loader — the `--mmproj` projector `LoadQwen3VLVisionFromClipMmproj`, which is the same Qwen3-VL tower read out of a second `clip` GGUF and which until then made `--language-model-only` zero every limit, refuse every image, and still pay for the projector. Construct-without-initialise, not skip-the-construction: geometry still resolves, the projector file is still opened and still refused by name if this build cannot use it, the loader's structural accounting does not move, and a skipped tower REFUSES BY NAME rather than reading empty buffers (`utils.py:693-705`). **RSS MEASURED 2026-08-24 for ONE of the three sites (#1358):** `scripts/mm/tower_skip_rss.sh --model-kind qwen3-vl` on `thor:gpu0` under an `rc` lease, at `41ab550b9`, `--device cpu`, `Qwen/Qwen3-VL-4B-Instruct` @ `ebb281ec` staged to worker-local disk — peak host RSS 10209501184 B default against 8553709568 B with `--language-model-only`, a saving of **1655791616 B = 1.542 GiB**, and 1655992320 B on the SWAPPED pair (spread 200704 B against a 192512 B leg-to-leg repeat; the two binaries were sha256-identical). That MEETS the 1495251763 B threshold declared before any number existed, on BOTH pairs, and it is 99.7% of the 1661390848 B tower predicted from the checkpoint's own headers. THREE caveats travel with it: about half the saving is #1359's bf16→host-f32 widening, whose Qwen3-VL half has since LANDED, so a rerun reads about 0.774 GiB rather than 1.542 and that halving is CORRECT rather than a regression — the flag now frees the tower the checkpoint ships instead of the tower plus our widening, and `muse-glimmer-30b` still widens (blocked on #2166); it is LOAD-TIME peak RSS, because `ForwardQwen3VLForConditionalGeneration` refuses text-only input through this arch and the arms therefore stop at `/health`; and only HALF 1 is asserted, the pre-L3 `edbc47ce0` within-2% half being a separate run that stays owed. **RERUN 2026-08-28 (#1359 verification), the CURRENT figure:** the same harness at `525d2b991` on `dgx:gpu0` under an `rc` lease read peak host RSS 9381281792 B default against 8554364928 B with `--language-model-only`, a saving of **826916864 B = 0.770 GiB**, and 826576896 B on the SWAPPED pair (spread 339968 B against a 24576 B leg-to-leg repeat; binaries again sha256-identical). That MEETS the live 747625881 B threshold — 90% of the 830695424 B tower the checkpoint ships, halved with #1359 exactly as the resident figure was — on BOTH pairs, and it is 99.5% of that tower. The predicted halving is therefore now MEASURED at **0.499x**, not expected. The same run VERIFIES #1359's Qwen3-VL half: the default arm recovered 828219392 B = 0.771 GiB (10209501184 → 9381281792), **99.7% of the 830695424 B predicted from the checkpoint headers**, while the tower-free `--language-model-only` CONTROL arm moved only **+655360 B = +0.0077%** against a 2% bound, which is what makes the attribution sound across the host change from `thor` to `dgx`. #1359 does NOT close on it: #2166 (muse-glimmer, blocked because its `kF32` gate fixture is a synthetic f32 LCG that never rounds through bf16, so regenerating that golden inside the change it gates would be circular) and #2173 (Gemma-4 vision, landed UNREACHED) both remain owed. STILL NOT CLAIMED: any RSS number for `muse-glimmer` (owed, needs ~56 G of worker-local disk) or for the `--mmproj` site, any GPU/VRAM figure, and any throughput figure. Evidence `docs/bench-evidence/tower-skip-rss-qwen3vl-thor-20260824{,.legs}.log` and `docs/bench-evidence/tower-skip-rss-qwen3vl-dgx-20260828{,.legs}.log`. **KERNEL GATE L4 RESOLVED 2026-08-19 as a TRACKED EXCEPTION (#607, #414):** we do NOT mirror the `text_only` conjunct of `qwen3_next.py:324-331`, because mirroring is not representable at our seams. Upstream conjoins it because its fused Triton kernel indexes `cos_sin_cache` by 1-D `positions` and cannot express MRoPE (`qwen3_next.py:323`, `# TODO: support MRoPE`), falling back to an eager arm whose `self.rotary_emb` IS the MRoPE module. Ours are not those two arms: `vt::AttnQkNormRopeGate` takes NO positions, only a precomputed per-token `cos_sin` cache that `qwen3_5.cpp::BuildMropeCosSinHost` fills with the interleaved 3-section MRoPE selection, so our FUSED arm is the MRoPE arm while our eager arm (`vt::RopeNeox` on 1-D positions) has no MRoPE spelling. Conjoining `text_only` would select 1-D RoPE on exactly the configuration the conjunct protects and would break the landed M3-b image and M3d video STRICT 32/32 gates. Argued in `specs/multimodal-track.md` §1.6. What the exception does NOT excuse is the DENOMINATOR: #414's defect is a benchmark configuration, and `scripts/dgx-online-serving.sh` still launched the oracle without `--language-model-only` while `tools/bench/run_serve_low.py` passed it, so the two harnesses disagreed about the oracle's own configuration and the next canonical campaign would have reproduced the flattered ratios the 2026-08-13 series superseded. Both now pass it and `scripts/check-oracle-denominator-flags.py` keeps them agreeing. NO product code path changed, so no token gate and no measurement is claimed and no published number is withdrawn. Owed and filed in flow: #1340 (`VT_FUSE_ATTN_PREAMBLE=0` on the MRoPE path substitutes 1-D RoPE instead of refusing; needs a GPU VL token gate) and #1345 (the three in-process `LLM(...)` bench harnesses leave `language_model_only` at False with no way to set it). | T1 | `vllm/multimodal/{inputs.py,hasher.py:50,processing/processor.py:1663,processing/inputs.py:62}`; `vllm/model_executor/models/qwen3_vl.py:{1400,1233}`; `vllm/v1/core/encoder_cache_manager.py:17`; transformers `image_processing_qwen2_vl.py:62`, `image_processing_backends.py:327`; tests `tests/multimodal/test_processing.py`, `tests/multimodal/test_hasher.py`, `tests/v1/core/test_encoder_cache_manager.py` @ `e24d1b24` | `src/vllm/multimodal/hasher.cpp`, `src/vllm/multimodal/qwen3vl_processor.cpp`, `include/vllm/multimodal/{inputs.h,hasher.h,qwen3vl_processor.h}`; `src/vllm/v1/core/encoder_cache_manager.cpp` + `include/vllm/v1/core/encoder_cache_manager.h`; additive inert fields `include/vllm/v1/request.h` + `src/vllm/v1/request.cpp` + `include/vllm/v1/engine/types.h`; `extra_keys` seam `include/vllm/v1/kv_offload/lmcache/chunked_token_database.h` + `.cpp`; M0 `scripts/mm/m0_oracle_capture.py`; L1 limits `include/vllm/config/multimodal.h` + `include/vllm/multimodal/processing/context.h` + `src/vllm/multimodal/processing/context.cpp`, refusal type relocated to `include/vllm/v1/engine/validation_error.h`; L2 flags+ABI+call site `src/vllm/config/multimodal.cpp` (`ParseLimitMmPerPromptJson`) + `src/vllm/entrypoints/openai/server_main.cpp` + `include/vllm.h` (ABI v19) + `src/capi/vllm_c.cpp` + `EngineParams::multimodal`/`LoadedEngine::mm_config()` + `src/vllm/entrypoints/openai/chat_mm.cpp` (`ChatPartModality`, `ValidateChatMmLimits`, `Qwen3VLChatSupportedMmLimits`) — anchors `src/vllm/multimodal/hasher.cpp:56`, `src/vllm/config/multimodal.cpp:49`, `src/vllm/entrypoints/openai/chat_mm.cpp:295,311`, `include/vllm.h:197,403`; L3 tower skip: the predicate `include/vllm/model_executor/models/interfaces.h::SkipTowerForModalities` + `src/vllm/model_executor/models/interfaces.cpp::SkipTowerForModalities`, the borrow `include/vllm/model_executor/models/model_registry.h::ModelSource` (`multimodal`) and the observable `include/vllm/model_executor/models/model_registry.h::skipped_towers` / `include/vllm/entrypoints/model_loader.h::skipped_towers`, the three gated loads `src/vllm/model_executor/models/qwen3_vl.cpp::LoadQwen3VLWeights`, `src/vllm/model_executor/models/muse_glimmer_weights.cpp::LoadMuseGlimmerForConditionalGenerationWeights` and `src/vllm/entrypoints/model_loader.cpp::LoadQwen3VLVisionFromClipMmproj`, the reporting sites `src/vllm/model_executor/models/qwen3_vl_registry.cpp::skipped_towers` / `src/vllm/model_executor/models/muse_glimmer_registry.cpp::skipped_towers` and the server line `src/vllm/entrypoints/openai/server_main.cpp::skipped_towers`; RSS procedure `scripts/mm/tower_skip_rss.sh`; L4 denominator gate `scripts/check-oracle-denominator-flags.py` + the `--language-model-only` oracle arm of `scripts/dgx-online-serving.sh` | `tests/vllm/multimodal/test_qwen3vl_processor.cpp` (processor-parity 23/23 BIT-identical vs the M0 oracle fixture `tests/vllm/multimodal/fixtures/qwen3vl/`, RED-first: wrong normalize shift -> 1.2M mismatches); `tests/vllm/v1/core/test_encoder_cache_manager.cpp` 32/32. Text-inertness: `test_request`/`test_engine_types`/`test_lmcache_codec`/`test_lmcache_key_agreement`/`test_openai_conformance` all green standalone; SACRED CUDA 27B/35B/Coder = GPU inertness proof; `check-device-leakage` OK — anchor `tests/vllm/multimodal/test_qwen3vl_processor.cpp:59`. L1 limits: `tests/vllm/config/test_multimodal_config.cpp` 7/7 (21 assertions) + `tests/vllm/multimodal/test_processing_limits.cpp` 19/19 (78 assertions), porting `tests/multimodal/test_processing.py:902-941,944-985`, `tests/entrypoints/multimodal/llm/test_mm_embeds_only.py:41-49` and `tests/entrypoints/unit_tests/test_chat_utils.py:1498-1560` @ `5559679229bc`; mutations proven RED: map-before-flag precedence, the dropped throw, the dropped `min()` fold. L2 (aarch64 `build-test-cpu-arm64` lane, `-DVLLM_CPP_CUDA=OFF`): `tests/vllm/entrypoints/openai/test_serve_mm_limits.cpp` 11/11 (109 assertions, flags + the parser's upstream refusals + the builtin-only reach of `extra="forbid"`) + `test_chat_mm` 11/11 (126) + `test_openai_api_server` 56/56 (CASES; its assertion count is timing-dependent — 632/648/651 across three runs of one binary, so only the case count is quotable — the HTTP 400 arm proven against BOTH a 500 and a truncated 200); RED-first behavioural: `CHECK(500 == 400)`, `CHECK("InternalServerError" == BadRequestError)` and `CHECK(200 == 400)`; mutations proven RED: flag→config plumbing dropped, the call-site wiring dropped, and the refusal re-typed off `InputValidationError` (which lands as the 500 L1's design avoided). L3: `tests/vllm/models/test_tower_skip.cpp` 6/6 (80 assertions — the decision, both tower loaders, text-path BIT-identity across the flag, the refuse-by-name placeholder, and the reachability case that alone reds when `source.multimodal = ¶ms.multimodal` is deleted) + `tests/vllm/entrypoints/test_gguf_mmproj_reach.cpp` 10/10 (32) for the `--mmproj` arm; mutations proven RED: deleting the `--mmproj` gate turns the zero-limit case red (4 assertions, the reader's own "inventing" refusal reappearing is the receipt that the read happened), and the reviewer's `(void)mm_config; vision_skipped=false; vision_loaded=false` — which left every suite green before the repair wave — now reds `test_tower_skip` by THROWING on the missing `model.visual.*`. NOT gated in CI: what `skipped_towers()` reports on the `--mmproj` arm (env-gated; a synthetic language GGUF carries no tokenizer, so no `LoadedEngine` is built from it) and every RSS number. L4: `tests/scripts/test_check_oracle_denominator_flags.py` 11/11, RED-first behavioural (the checker on the pre-L4 tree exits 1 naming `dgx-online-serving.sh:487` and `:498` of 3 discovered launches) and mutation-proven (removing the flag from the canonical driver in a scratch copy returns exit 1 with exactly one violation, the exempt q3mxfp4 arm staying exempt) | [multimodal-track.md](specs/multimodal-track.md) §3 (M0/M1) | `READY` | - | +| `ENG-MM-INPUT-PIPELINE` | **Multimodal INPUT pipeline + encoder-cache engine seam (M1), INERT when no mm input.** The C++ mirror of `vllm/multimodal/`: `MultiModalKwargs`/`MultiModalFeatureSpec`/`MultiModalInputs`, the `MultiModalHasher` mm-hash (blake3), the Qwen3-VL image processor (smart_resize + fused rescale/normalize + patchify -> `pixel_values`+`image_grid_thw`) and placeholder-token expansion, plus the `EncoderCacheManager` (+`ComputeMmEncoderBudget`) and the LMCache `extra_keys` seam. Additive `mm_features` carried on `Request`/`EngineCoreRequest`; with NO mm input every field is empty and every path is byte-identical to the text engine. Processor output is BIT/BYTE-identical to the vLLM 0.25.0 oracle (M0 fixture). Does NOT build the vision tower / embed-merge (M2). **SERVING wiring (ROAD-V1-MM `MM-SERVE-ENGINE`, 2026-07-28, `CLAIM-MM-SERVING-W2`):** the OpenAI server now carries the parsed `MultiModalInputs` into the engine — additive `LLMEngine`/`AsyncLLM` `add_request(MultiModalInputs)`+`generate(MultiModalInputs)` overloads via `InputProcessor::process_inputs_mm` (mirror `input_processor.py:333-379`, empty mm_features == the tokens path), the chat-template placeholder-STRING helpers (`get_placeholder_str`/`_add_placeholder` mirror), and the serving_chat `MultiModalChatFn` seam (default unset ⇒ text byte-identical). **SEAM BODY (ROAD-V1-MM `MM-SERVE-E2E` W3, 2026-07-28, `CLAIM-MM-SERVING-E2E`):** `MakeQwen3VLImageChatFn` (chat_mm.cpp) is the seam body the server sets — messages → marker-inject → chat template → `EncodeWithSpecialTokens` (the single image_pad marker → one image_token_id) → `RouteImageRgb` EXPAND to 196 image tokens + mm_features; wired in `examples/server/main.cpp` (guarded on `preprocessor_config.json`; text-only unset ⇒ byte-identical). Gated `test_chat_mm` 8/8 + `test_openai_serving` (seam invoked + routed). **ENGINE MM-FORWARD LANDED (ROAD-V1-MM `MM-SERVE-E2E`, 2026-07-28, `CLAIM-ENGINE-MM-FORWARD`):** the engine model runner now HAS an mm forward — `ModelForwardInput` gains an ADDITIVE default-nullopt `std::optional mm` (merged inputs_embeds + 3-D MRoPE positions + DeepStack, borrowed handles; nullopt-for-text ⇒ shared runner path byte-identical BY CONSTRUCTION), `Qwen3VLForConditionalGeneration` is `REGISTER_VLLM_MODEL`-registered (`qwen3_vl_registry.cpp`), and the registered forward FOLDS the M2c decode into `ModelRegistry::Forward` via the SHARED `Qwen3VLForwardStepLastLogits` (`Qwen3VLGenerateGreedyViaRegistry` drives every step through the registry). GPU token-exact gate `test_qwen3vl_registry_e2e` (image→text THROUGH `ModelRegistry::Forward` == M2c golden 32/32 STRICT, dgx.casa GB10); text inertness `test_runner` 16/16 + `test_scheduler` 36/36 + `test_model_registry` 24/24 + `test_chat_mm` 8/8 + `test_openai_serving` 41/41 all green. RESIDUAL: the FULL in-runner scheduler-fed tower run (batched-loop mm building the field from staged encoder outputs) + the real server `/v1/chat/completions` GPU e2e — recipe in `specs/mm-serving.md`. **INPUT LIMITS L1 LANDED (#607, 2026-08-13):** the per-modality `limit_per_prompt` + `GetLimitPerPrompt` precedence (`language_model_only` ⇒ 0 BEFORE the map, else the map, else 999) and the refusal that gives those numbers effect — `AllowedMmLimits` folding by `min()` against the model's own ceiling, `ValidateNumItems` with upstream's exact message, and both call sites with the `enable_mm_embeds` escape. NO serve surface and NO live call site: nothing constructs a `MultiModalConfig` on a request yet, which is L2's. **INPUT LIMITS L2 LANDED (#607, #686, 2026-08-14):** the flags (`--[no-]language-model-only`, `--limit-mm-per-prompt ''`; `arg_utils.py:555-556,1276-1279,1691-1692` over `ParseLimitMmPerPromptJson`, the port of `multimodal.py:212-236` + the DummyOptions dataclasses `:17-43`), the C-ABI fields (`vllm_model_params.language_model_only`/`.limit_mm_per_prompt`, **ABI v19**), and the LIVE CALL SITE: `ValidateChatMmLimits` (`chat_utils.py:648-662`) runs as step 0 of `MakeQwen3VLImageChatFn` over a `BaseProcessingInfo` folding `LoadedEngine::mm_config()` with `Qwen3VLChatSupportedMmLimits() == {"image": 1}` — the seam's own ceiling, which is the `min()` fold operand #686 recorded as undeclared. A three-image request is now HTTP 400 with upstream's message rather than an opaque 500 / a truncated answer. NOT claimed AT L2: any memory win — L3 gates tower construction, and its saving is MEASURED only for `Qwen3-VL-4B-Instruct` (host RSS at load, not VRAM, and half 1 of a two-half gate); `muse-glimmer` and the `--mmproj` site stay unmeasured. Still unwired: the `process_inputs_mm` call site (`context.py:461`), blocked on the per-model `get_supported_mm_limits()` hook. **TOWER SKIP L3 (#607):** the predicate `SkipTowerForModalities` (the port of `interfaces.py:293` — `all(...) == 0` over the tower's OWN modality set, ALL not ANY, reached through `GetLimitPerPrompt` so `--limit-mm-per-prompt '{"image":0,"video":0}'` skips exactly as `--language-model-only` does), the borrow that carries the engine's limits into a load (`ModelSource::multimodal`), and the observable that says what was freed (`LoadedModel::skipped_towers` / `LoadedEngine::skipped_towers`, printed once by the server). All THREE production tower loads are gated: `LoadQwen3VLWeights`, `LoadMuseGlimmerForConditionalGenerationWeights`, and — added in the L3 repair wave, missed by the first cut because it lives in the entrypoint rather than in a model loader — the `--mmproj` projector `LoadQwen3VLVisionFromClipMmproj`, which is the same Qwen3-VL tower read out of a second `clip` GGUF and which until then made `--language-model-only` zero every limit, refuse every image, and still pay for the projector. Construct-without-initialise, not skip-the-construction: geometry still resolves, the projector file is still opened and still refused by name if this build cannot use it, the loader's structural accounting does not move, and a skipped tower REFUSES BY NAME rather than reading empty buffers (`utils.py:693-705`). **RSS MEASURED 2026-08-24 for ONE of the three sites (#1358):** `scripts/mm/tower_skip_rss.sh --model-kind qwen3-vl` on `thor:gpu0` under an `rc` lease, at `41ab550b9`, `--device cpu`, `Qwen/Qwen3-VL-4B-Instruct` @ `ebb281ec` staged to worker-local disk — peak host RSS 10209501184 B default against 8553709568 B with `--language-model-only`, a saving of **1655791616 B = 1.542 GiB**, and 1655992320 B on the SWAPPED pair (spread 200704 B against a 192512 B leg-to-leg repeat; the two binaries were sha256-identical). That MEETS the 1495251763 B threshold declared before any number existed, on BOTH pairs, and it is 99.7% of the 1661390848 B tower predicted from the checkpoint's own headers. THREE caveats travel with it: about half the saving is #1359's bf16→host-f32 widening, whose Qwen3-VL half has since LANDED, so a rerun should read about 0.774 GiB rather than 1.542 and that halving is CORRECT rather than a regression — the flag now frees the tower the checkpoint ships instead of the tower plus our widening, and `muse-glimmer-30b` still widens (blocked on #2166); it is LOAD-TIME peak RSS, because `ForwardQwen3VLForConditionalGeneration` refuses text-only input through this arch and the arms therefore stop at `/health`; and only HALF 1 is asserted, the pre-L3 `edbc47ce0` within-2% half being a separate run that stays owed. STILL NOT CLAIMED: any RSS number for `muse-glimmer` (owed, needs ~56 G of worker-local disk) or for the `--mmproj` site, any GPU/VRAM figure, and any throughput figure. Evidence `docs/bench-evidence/tower-skip-rss-qwen3vl-thor-20260824{,.legs}.log`. **KERNEL GATE L4 RESOLVED 2026-08-19 as a TRACKED EXCEPTION (#607, #414):** we do NOT mirror the `text_only` conjunct of `qwen3_next.py:324-331`, because mirroring is not representable at our seams. Upstream conjoins it because its fused Triton kernel indexes `cos_sin_cache` by 1-D `positions` and cannot express MRoPE (`qwen3_next.py:323`, `# TODO: support MRoPE`), falling back to an eager arm whose `self.rotary_emb` IS the MRoPE module. Ours are not those two arms: `vt::AttnQkNormRopeGate` takes NO positions, only a precomputed per-token `cos_sin` cache that `qwen3_5.cpp::BuildMropeCosSinHost` fills with the interleaved 3-section MRoPE selection, so our FUSED arm is the MRoPE arm while our eager arm (`vt::RopeNeox` on 1-D positions) has no MRoPE spelling. Conjoining `text_only` would select 1-D RoPE on exactly the configuration the conjunct protects and would break the landed M3-b image and M3d video STRICT 32/32 gates. Argued in `specs/multimodal-track.md` §1.6. What the exception does NOT excuse is the DENOMINATOR: #414's defect is a benchmark configuration, and `scripts/dgx-online-serving.sh` still launched the oracle without `--language-model-only` while `tools/bench/run_serve_low.py` passed it, so the two harnesses disagreed about the oracle's own configuration and the next canonical campaign would have reproduced the flattered ratios the 2026-08-13 series superseded. Both now pass it and `scripts/check-oracle-denominator-flags.py` keeps them agreeing. NO product code path changed, so no token gate and no measurement is claimed and no published number is withdrawn. Owed and filed in flow: #1340 (`VT_FUSE_ATTN_PREAMBLE=0` on the MRoPE path substitutes 1-D RoPE instead of refusing; needs a GPU VL token gate) and #1345 (the three in-process `LLM(...)` bench harnesses leave `language_model_only` at False with no way to set it). | T1 | `vllm/multimodal/{inputs.py,hasher.py:50,processing/processor.py:1663,processing/inputs.py:62}`; `vllm/model_executor/models/qwen3_vl.py:{1400,1233}`; `vllm/v1/core/encoder_cache_manager.py:17`; transformers `image_processing_qwen2_vl.py:62`, `image_processing_backends.py:327`; tests `tests/multimodal/test_processing.py`, `tests/multimodal/test_hasher.py`, `tests/v1/core/test_encoder_cache_manager.py` @ `e24d1b24` | `src/vllm/multimodal/hasher.cpp`, `src/vllm/multimodal/qwen3vl_processor.cpp`, `include/vllm/multimodal/{inputs.h,hasher.h,qwen3vl_processor.h}`; `src/vllm/v1/core/encoder_cache_manager.cpp` + `include/vllm/v1/core/encoder_cache_manager.h`; additive inert fields `include/vllm/v1/request.h` + `src/vllm/v1/request.cpp` + `include/vllm/v1/engine/types.h`; `extra_keys` seam `include/vllm/v1/kv_offload/lmcache/chunked_token_database.h` + `.cpp`; M0 `scripts/mm/m0_oracle_capture.py`; L1 limits `include/vllm/config/multimodal.h` + `include/vllm/multimodal/processing/context.h` + `src/vllm/multimodal/processing/context.cpp`, refusal type relocated to `include/vllm/v1/engine/validation_error.h`; L2 flags+ABI+call site `src/vllm/config/multimodal.cpp` (`ParseLimitMmPerPromptJson`) + `src/vllm/entrypoints/openai/server_main.cpp` + `include/vllm.h` (ABI v19) + `src/capi/vllm_c.cpp` + `EngineParams::multimodal`/`LoadedEngine::mm_config()` + `src/vllm/entrypoints/openai/chat_mm.cpp` (`ChatPartModality`, `ValidateChatMmLimits`, `Qwen3VLChatSupportedMmLimits`) — anchors `src/vllm/multimodal/hasher.cpp:56`, `src/vllm/config/multimodal.cpp:49`, `src/vllm/entrypoints/openai/chat_mm.cpp:295,311`, `include/vllm.h:197,403`; L3 tower skip: the predicate `include/vllm/model_executor/models/interfaces.h::SkipTowerForModalities` + `src/vllm/model_executor/models/interfaces.cpp::SkipTowerForModalities`, the borrow `include/vllm/model_executor/models/model_registry.h::ModelSource` (`multimodal`) and the observable `include/vllm/model_executor/models/model_registry.h::skipped_towers` / `include/vllm/entrypoints/model_loader.h::skipped_towers`, the three gated loads `src/vllm/model_executor/models/qwen3_vl.cpp::LoadQwen3VLWeights`, `src/vllm/model_executor/models/muse_glimmer_weights.cpp::LoadMuseGlimmerForConditionalGenerationWeights` and `src/vllm/entrypoints/model_loader.cpp::LoadQwen3VLVisionFromClipMmproj`, the reporting sites `src/vllm/model_executor/models/qwen3_vl_registry.cpp::skipped_towers` / `src/vllm/model_executor/models/muse_glimmer_registry.cpp::skipped_towers` and the server line `src/vllm/entrypoints/openai/server_main.cpp::skipped_towers`; RSS procedure `scripts/mm/tower_skip_rss.sh`; L4 denominator gate `scripts/check-oracle-denominator-flags.py` + the `--language-model-only` oracle arm of `scripts/dgx-online-serving.sh` | `tests/vllm/multimodal/test_qwen3vl_processor.cpp` (processor-parity 23/23 BIT-identical vs the M0 oracle fixture `tests/vllm/multimodal/fixtures/qwen3vl/`, RED-first: wrong normalize shift -> 1.2M mismatches); `tests/vllm/v1/core/test_encoder_cache_manager.cpp` 32/32. Text-inertness: `test_request`/`test_engine_types`/`test_lmcache_codec`/`test_lmcache_key_agreement`/`test_openai_conformance` all green standalone; SACRED CUDA 27B/35B/Coder = GPU inertness proof; `check-device-leakage` OK — anchor `tests/vllm/multimodal/test_qwen3vl_processor.cpp:59`. L1 limits: `tests/vllm/config/test_multimodal_config.cpp` 7/7 (21 assertions) + `tests/vllm/multimodal/test_processing_limits.cpp` 19/19 (78 assertions), porting `tests/multimodal/test_processing.py:902-941,944-985`, `tests/entrypoints/multimodal/llm/test_mm_embeds_only.py:41-49` and `tests/entrypoints/unit_tests/test_chat_utils.py:1498-1560` @ `5559679229bc`; mutations proven RED: map-before-flag precedence, the dropped throw, the dropped `min()` fold. L2 (aarch64 `build-test-cpu-arm64` lane, `-DVLLM_CPP_CUDA=OFF`): `tests/vllm/entrypoints/openai/test_serve_mm_limits.cpp` 11/11 (109 assertions, flags + the parser's upstream refusals + the builtin-only reach of `extra="forbid"`) + `test_chat_mm` 11/11 (126) + `test_openai_api_server` 56/56 (CASES; its assertion count is timing-dependent — 632/648/651 across three runs of one binary, so only the case count is quotable — the HTTP 400 arm proven against BOTH a 500 and a truncated 200); RED-first behavioural: `CHECK(500 == 400)`, `CHECK("InternalServerError" == BadRequestError)` and `CHECK(200 == 400)`; mutations proven RED: flag→config plumbing dropped, the call-site wiring dropped, and the refusal re-typed off `InputValidationError` (which lands as the 500 L1's design avoided). L3: `tests/vllm/models/test_tower_skip.cpp` 6/6 (80 assertions — the decision, both tower loaders, text-path BIT-identity across the flag, the refuse-by-name placeholder, and the reachability case that alone reds when `source.multimodal = ¶ms.multimodal` is deleted) + `tests/vllm/entrypoints/test_gguf_mmproj_reach.cpp` 10/10 (32) for the `--mmproj` arm; mutations proven RED: deleting the `--mmproj` gate turns the zero-limit case red (4 assertions, the reader's own "inventing" refusal reappearing is the receipt that the read happened), and the reviewer's `(void)mm_config; vision_skipped=false; vision_loaded=false` — which left every suite green before the repair wave — now reds `test_tower_skip` by THROWING on the missing `model.visual.*`. NOT gated in CI: what `skipped_towers()` reports on the `--mmproj` arm (env-gated; a synthetic language GGUF carries no tokenizer, so no `LoadedEngine` is built from it) and every RSS number. L4: `tests/scripts/test_check_oracle_denominator_flags.py` 11/11, RED-first behavioural (the checker on the pre-L4 tree exits 1 naming `dgx-online-serving.sh:487` and `:498` of 3 discovered launches) and mutation-proven (removing the flag from the canonical driver in a scratch copy returns exit 1 with exactly one violation, the exempt q3mxfp4 arm staying exempt) | [multimodal-track.md](specs/multimodal-track.md) §3 (M0/M1) | `READY` | - | | `ENG-MM-VISION-TOWER` | **Qwen3-VL vision tower `Qwen3_VisionTransformer` (M2a), proven faithful vs vLLM 0.25.0 in isolation.** The reusable vision half of the whole Qwen3-VL family + Qwen3.6 (27B/35B share this exact tower). Pure-additive C++ forward composed from public vt:: ops: patch-embed (Conv3d-as-matmul + bias), host pos-embed bilinear-interp+spatial-merge-reorder, 24 ViT blocks (LayerNorm + vision attention with partial-rotary NeoX vision RoPE via `vt::RopeFromCache` + non-causal `vt::Attention(causal=false)` + tanh-GELU MLP), patch merger (LayerNorm + exact-erf-GELU + 2 FCs), DeepStack 3 post-shuffle-norm mergers at layers 5/11/17 → `[196,10240]`. Adds 2 additive elementwise vt ops (`GeluTanh`/`GeluErf`). NO runner/model/registry edit → text engines byte-identical by construction. Proven faithful in ISOLATION; the merge into `input_embeds` + the MRoPE/DeepStack text backbone + the e2e image gate are M2b/M2c. | T1 | `vllm/model_executor/models/qwen3_vl.py` `Qwen3_VisionPatchEmbed:347`, `Qwen3_VisionBlock:413`, `Qwen3_VisionPatchMerger:467`, `Qwen3_VisionTransformer:519`, `forward:800`, `pos_embed_interpolate_native:277`, `rot_pos_emb:667`; `qwen2_5_vl.py::Qwen2_5_VisionAttention.forward:397`; `rotary_embedding/common.py::ApplyRotaryEmb.forward_static:151` @ `e24d1b24` | `src/vllm/model_executor/models/qwen3_vl_vision.{h,cpp}`; 2 vt ops `include/vt/ops.h` + `src/vt/ops.cpp` + `src/vt/cuda/cuda_layernorm.cu` + `src/vt/cpu/cpu_layernorm.cpp`; dumps `scripts/mm/m2a_tower_{ref,weight}_dump.py`; fixtures `tests/vllm/multimodal/fixtures/qwen3vl_tower/` | `tests/vllm/multimodal/test_qwen3vl_tower.cpp` — 4 RED-first tower gates vs the dumped vLLM-0.25.0 reference 348/348 (patch-embed 2.1e-3, block0 6.8e-3, merger 6.5e-2, DeepStack 1.2e-2/3.3e-2/4.4e-2, full tower 5.1e-2; pos-embed 2.5e-3 + rope 1.9e-3 TIGHT); bf16-depth envelope RCA'd; RED = rope disabled → block0 0.149/tower 0.75/6 fails; cutlass-ON+FA2 banner; clean `-Werror`; compute-sanitizer 0 — anchor `tests/vllm/multimodal/test_qwen3vl_tower.cpp:96` | [multimodal-track.md](specs/multimodal-track.md) §3 (M2a) | `ACTIVE` | `CLAIM-MULTIMODAL-M2A` | | `ENG-MM-TEXT-BACKBONE` | **Qwen3-VL text-backbone numeric contracts `Qwen3VLGetRopeIndex`/`Qwen3VLMergeMultimodal`/`Qwen3VLComputeDeepstack` (M2b/M2c), unit-green vs vLLM 0.25.0.** The deterministic pieces that fork the plain Qwen3-dense text path for a vision-conditioned decode: (1) MRoPE 3-D `get_rope_index` positions [3,T] (image tokens get (t,h,w) grid positions, text sequential); (2) the 3-section MRoPE APPLICATION — proven to be the EXISTING `vt::RopeFromCache` mrope path (positions [3,T] + `mrope_section=[24,20,20]` interleaved), faithful to `MRotaryEmbedding.forward_native` for Qwen3-VL's exact config; (3) `_compute_deepstack_embeds` scatter → [L,T,H] decoder-injection tensor; (4) `_merge_multimodal_embeddings` masked scatter of the tower's `[:,:2560]` into `input_embeds`. Pure-additive TU — NO shared dense forward / runner / registry edit → text engines byte-identical by construction. The e2e image forward (VL weight loader + forked MRoPE/DeepStack decode loop) is the remaining M2c wire-up. | T1 | `vllm/model_executor/models/qwen3_vl.py` `_get_mrope_input_positions:2567`, `_iter_mm_grid_hw:2482`, `_compute_deepstack_embeds:2761`, `Qwen3LLMModel.forward` deepstack `:1589`; `vllm/model_executor/models/utils.py::_merge_multimodal_embeddings:524`; `vllm/model_executor/layers/rotary_embedding/mrope.py` MRotaryEmbedding @ `e24d1b24` | `src/vllm/model_executor/models/qwen3_vl_text.{h,cpp}`; existing `vt::RopeFromCache` mrope path (`src/vt/{cpu,cuda}/*`); dump `scripts/mm/m2b_text_ref_dump.py`; fixtures `tests/vllm/multimodal/fixtures/qwen3vl_text/` — anchor `src/vllm/model_executor/models/qwen3_vl_text.cpp:9` | `tests/vllm/multimodal/test_qwen3vl_text.cpp` — 4 RED-first gates vs the dumped vLLM-0.25.0 reference 85/85 (get_rope_index BIT-exact [3,204], delta −182; MRoPE q rel-L2 1.5e-3 / k 1.5e-3, RED interleaved-off >5e-2; DeepStack + merge BIT-exact); CPU-only, no weights; clean CPU `-Werror` — anchor `tests/vllm/multimodal/test_qwen3vl_text.cpp:99` | [multimodal-track.md](specs/multimodal-track.md) §3 (M2b/M2c) | `ACTIVE` | `CLAIM-MULTIMODAL-M2BC` | | `ENG-MM-QWEN36-VL-FORWARD` | **Qwen3.6-27B (`Qwen3_5ForConditionalGeneration`) GDN-hybrid VL forward — IMAGE (M3-b) + VIDEO (M3d) BOTH e2e, STRICT gates PASS 32/32. Our own gate model's image+video paths now work end-to-end (speed pending).** The genuinely-new integration completing our own gate model's mm paths: fork the landed bf16 `Qwen3_5DenseModel` GDN-hybrid forward (48 GDN + 16 full-attn) on gated, default-off points so a text-only 27B request stays byte-identical — (a) `inputs_embeds` entry (embed ids + `Qwen3VLMergeMultimodal` scatter of the 27B tower merger `[N,5120]` into the visual-token rows; 27B has EMPTY `deepstack_visual_indexes` ⇒ NO DeepStack); (b) 3-section MRoPE (`mrope_section=[11,11,10]` interleaved, rotary_dim 64, theta 1e7) in the 16 full-attn layers only via the proven `vt::RopeFromCache` mrope path (GDN layers carry no rope); (c) mixed load = the M2a `Qwen3_VisionTransformer` (27B vision config, empty deepstack) bf16 tower + the bf16 GDN-hybrid LLM via the EXISTING `LoadQwen3_5Dense`. **M3d (2026-07-25) added VIDEO by REUSE:** the M3-b image driver refactored into a shared `VLGenerateCoreGdn`, image+video wrappers differ ONLY in the merge mask (`image_token` vs `video_token` across frames) + the get_rope_index (`Qwen3VLGetRopeIndex` vs `Qwen3VLGetRopeIndexVideo`); the M3c processor/windowed-tower/video-MRoPE are reused verbatim. | T1 | `vllm/model_executor/models/qwen3_5.py:389` (`Qwen3_5ForConditionalGeneration` subclasses `Qwen3VLForConditionalGeneration`; `visual = Qwen3_VisionTransformer`, modalities {"image","video"}); `qwen3_vl.py` `_process_video_input:2165`, `_get_mrope_input_positions:2567` video branch, `get_video_repl:1479`; the 27B `config.json` (`mrope_section=[11,11,10]`, empty `deepstack_visual_indexes`) @ `e24d1b24` / vLLM 0.25.0 | **M3-b + M3d BUILT + GATED 2026-07-25:** vision-only loader `LoadQwen3VLVisionWeights` (`src/vllm/model_executor/models/qwen3_vl.cpp`, 27B config) + shared `VLGenerateCoreGdn` + image driver `Qwen3_5VLGenerateGreedy` + **video driver `Qwen3_5VLGenerateGreedyVideo`** + `BuildMropeCosSinHost` + the `mrope_cos_sin` param on `DenseForwardLayers` (`src/vllm/model_executor/models/qwen3_5.cpp`, nullptr on every text caller ⇒ byte-identical; the video driver is purely additive, the shared text forward UNTOUCHED per `git diff --stat`) reusing M2a tower + `LoadQwen3_5Dense` bf16 LLM | **IMAGE:** golden `tests/vllm/multimodal/fixtures/qwen3_5_27b/` (STRICT sha256 `ead4b484…`); STRICT image gate PASS **32/32** (`test_qwen3_5_vl_e2e.cpp`, 54/54, re-run post-refactor). **VIDEO (M3d):** oracle `scripts/mm/m3d_video_oracle_capture.py` on the M3c synthetic clip (raw sha `8a111599…`, grid `[4,8,8]`, 64 video tokens) K=5 DETERMINISTIC ⇒ STRICT golden; **STRICT video gate PASS 32/32** (`test_qwen3_5_vl_video_e2e.cpp`, 27/27; near-tie gaps 0.0000 nats everywhere), fixtures `tests/vllm/multimodal/fixtures/qwen3_5_27b_video/`. Text-inertness 27B 235/235, 35B 315/315, Coder 138/138 (by construction); clean `-Werror` 0 warn; compute-sanitizer 0 on the 27B video forward. **SPEED MEASURED (2026-07-26, `CLAIM-MULTIMODAL-SPEED`): image c1 vs vLLM 0.25.0 GRAPHED — decode TPOT 225.0 ms/tok vs 226.9 = AT PARITY (0.99×), LLM prefill 326 ms vs vLLM TTFT 321 ms = at parity; vision tower WAS 2114 ms vs vLLM encode ≤~250 ms = ~10× (THE gap). TOWER LEVER EXECUTED (2026-07-26, `CLAIM-MULTIMODAL-SPEED-TOWER`, [multimodal-speed.md](specs/multimodal-speed.md) §7): nsys `cuda_gpu_kern_sum` attributed 98.9 % of the tower forward to the naive `vt::cuda::AttentionKernel` (56 ms/block; NOT QKV/FA2-routing); fixed by a warp-scoped online-softmax op `AttentionDenseFast` (separate op ⇒ `kAttention`/text byte-identical) + one-time resident-weight load ⇒ per-image tower 2114 → 148 ms (14.3×), **0.59× vs vLLM eager encode = FASTER**. STRICT image/video e2e HELD 32/32 (+4B DeepStack 32/32), `test_ops_attention` 37239/37239, 27B text SACRED 235/235, compute-sanitizer memcheck 0, clean `-Werror`. `benchmark_binding=false`, single-seq driver (no c2+/server). Remaining: batched/graphed mm serving (c2+) + audio our-side — DONE bar not yet met.** | [multimodal-track.md](specs/multimodal-track.md) §M3 + [multimodal-speed.md](specs/multimodal-speed.md) §7 + §8 (decode lever #2 CLOSED 2026-07-27: on-GPU greedy argmax + decode embed round-trip removed on `VLGenerateCoreGdn`; bit-exact — image/video STRICT 32/32 held; 27B decode NEUTRAL at the ~222 ms bandwidth floor) + §9 (lever #3 FIRST BRICK 2026-07-27, `CLAIM-MULTIMODAL-SPEED-GRAPH`: the shared `VLGenerateCoreGdn` decode step now routes through the production `Qwen3_5DenseDecodeGraph` cold→warm→replay captured decode — the mm decode is now GRAPH-CAPTURABLE, closing the un-graphed-eager-loop structural gap; S==B==1 bit-identical rebuild; token-exact HELD image/video STRICT 32/32 with 30 graph replays confirmed; A/B graphed 232.5 vs eager 233.4 ms/tok = NEUTRAL at the 27B bandwidth floor; the launch-overhead win + batched c2+ + serving ingestion are the recorded W-plan W1-W3) + §16 (vision-forward flash kernel 2026-07-28, `CLAIM-MM-SPEED-QWEN-IMAGE`: ATTRIBUTION-FIRST nsys attributed ~85% of the 148 ms tower forward to the dense attention `AttentionWarpKernel` [4.66 ms/block×27]; routed it to the §14 flash-tiled `vt::AttentionDenseFlash` [head_dim 72, byte-identical — per-warp math verbatim, only K/V from shared-mem tiles]. STRICT image/video e2e HELD 32/32 [27B+4B], `test_ops_attention` 37239/37239, goldens md5 UNCHANGED, nsys proof `AttentionDenseFlashKernel` 24 inst/zero warp, RED 30/46→46/46, sanitizer 0. A/B warp 148.3→flash 142.3 ms = 1.04× — the profile REFUTED a big lever: at t=784 the vision attention is serial-latency-bound not bandwidth-bound [audio §14 was 1.82× at t=1500], flash recovers only ~6 ms. **HONEST: the tower ALREADY BEATS vLLM — 142 ms vs ~250 ms eager encode = 0.57×**; image/video mm-forward is correctness-DONE + speed-BEATS-vLLM; residual = tensor-core MMA hd-72 attention [not needed for parity] + batched c2+/serving) | `ACTIVE` | `CLAIM-MULTIMODAL-SPEED-TOWER` + `CLAIM-MULTIMODAL-SPEED-DECODE` + `CLAIM-MULTIMODAL-SPEED-GRAPH` + `CLAIM-MM-SPEED-QWEN-IMAGE` | @@ -82,7 +82,7 @@ forensics: roadmap_v1.md and the parity ledger. | `ENG-MM-AUDIO-PIPELINE` | **AUDIO INPUT pipeline (audio-track A0+A1), the genuinely-new AUDIO modality on the modality-agnostic mm spine; INERT when no audio input.** Stands audio up on the smallest oracle-runnable vehicle `openai/whisper-small` (native `WhisperEncoder`; transformers 5.13.1 constructs it — unlike Gemma-4 which is oracle-blocked). The C++ Whisper-class audio processor `WhisperAudioProcessor`: canonical PCM16-mono WAV decode (`int16/32768.0`), identity resample at 16 kHz (genuine windowed-sinc DEFERRED, mirrors the image SmartResize/bicubic deferral), log-mel `input_features` `[80,3000]` (pad/truncate 480000 → torch.stft-equiv: reflect-pad `n_fft/2`, periodic Hann, hop 160, drop last frame, direct DFT over 201 bins → `abs(stft)^2` → `mel_filters.T@mag` → `log10(clamp 1e-10)` → `max(x,x.max()-8)` → `(x+4)/4`), audio placeholder expansion (`[0]`→`[0]*1500`, num_audio_tokens = `max_source_positions` = encoder output length), and `MultiModalHasher::HashAudioF32` (float32 1-D ndarray `"