diff --git a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD.md b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD.md index 0776d6485d..9159f014f9 100644 --- a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD.md +++ b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | Owns the Tenstorrent host-free forward per `.agents/specs/tenstorrent-host-free-forward.md` (issues #1604, #1625, #1627). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` (`ACTIVE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | Owns the Tenstorrent host-free forward per `.agents/specs/tenstorrent-host-free-forward.md` (issues #1604, #1625, #1627). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-MISTRAL.md b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-MISTRAL.md index 5649b70fa1..43d8b4efbf 100644 --- a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-MISTRAL.md +++ b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-MISTRAL.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-BACKEND-TENSTORRENT-MISTRAL` | `BACKEND-TENSTORRENT-MISTRAL` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-MISTRAL` | Owns the Tenstorrent Mistral port per `.agents/specs/tenstorrent-mistral.md`. Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-BACKEND-TENSTORRENT-MISTRAL` | `BACKEND-TENSTORRENT-MISTRAL` (`ACTIVE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-MISTRAL` | Owns the Tenstorrent Mistral port per `.agents/specs/tenstorrent-mistral.md`. Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN.md b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN.md index e0688b0b6b..b8f657eea1 100644 --- a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN.md +++ b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | `BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | Owns the Tenstorrent residual-golden spike per `.agents/specs/tenstorrent-residual-golden.md` (issue #289). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `SPIKE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | `BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` (`SPIKE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | Owns the Tenstorrent residual-golden spike per `.agents/specs/tenstorrent-residual-golden.md` (issue #289). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `SPIKE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER.md b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER.md index c081b6930b..84e402080d 100644 --- a/.agents/claims/CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER.md +++ b/.agents/claims/CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER` | `BACKEND-TENSTORRENT-TRACE-RUNNER` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-TRACE-RUNNER` | Owns the Tenstorrent trace-runner spike per `.agents/specs/tenstorrent-trace-runner.md` (issue #354). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `SPIKE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER` | `BACKEND-TENSTORRENT-TRACE-RUNNER` (`SPIKE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/BACKEND-TENSTORRENT-TRACE-RUNNER` | Owns the Tenstorrent trace-runner spike per `.agents/specs/tenstorrent-trace-runner.md` (issue #354). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `SPIKE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/claims/CLAIM-ENG-RELEASE-WINDOWS.md b/.agents/claims/CLAIM-ENG-RELEASE-WINDOWS.md index 7f9c6e48b7..586f58a774 100644 --- a/.agents/claims/CLAIM-ENG-RELEASE-WINDOWS.md +++ b/.agents/claims/CLAIM-ENG-RELEASE-WINDOWS.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-ENG-RELEASE-WINDOWS` | `ENG-RELEASE-WINDOWS` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/ENG-RELEASE-WINDOWS` | Owns the native Windows x86_64 pre-alpha release extension per `.agents/specs/windows-binary-release.md` (issue #117). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-ENG-RELEASE-WINDOWS` | `ENG-RELEASE-WINDOWS` (`ACTIVE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/ENG-RELEASE-WINDOWS` | Owns the native Windows x86_64 pre-alpha release extension per `.agents/specs/windows-binary-release.md` (issue #117). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/claims/CLAIM-LORA-RUNTIME-W2.md b/.agents/claims/CLAIM-LORA-RUNTIME-W2.md index 1cc6a4f51a..6ff8fa11b2 100644 --- a/.agents/claims/CLAIM-LORA-RUNTIME-W2.md +++ b/.agents/claims/CLAIM-LORA-RUNTIME-W2.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-LORA-RUNTIME-W2` | `LORA-RUNTIME` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/LORA-RUNTIME` | Owns the W2 packed + merged/TP LoRA layers per `.agents/specs/lora-adapter.md` (issue #278). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-LORA-RUNTIME-W2` | `LORA-RUNTIME` (`ACTIVE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/LORA-RUNTIME` | Owns the W2 packed + merged/TP LoRA layers per `.agents/specs/lora-adapter.md` (issue #278). Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/claims/CLAIM-WEIGHT-OFFLOAD-W0A.md b/.agents/claims/CLAIM-WEIGHT-OFFLOAD-W0A.md index b81cb2253a..57effa41ef 100644 --- a/.agents/claims/CLAIM-WEIGHT-OFFLOAD-W0A.md +++ b/.agents/claims/CLAIM-WEIGHT-OFFLOAD-W0A.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-WEIGHT-OFFLOAD-W0A` | `ENG-WEIGHT-OFFLOAD` | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/ENG-WEIGHT-OFFLOAD` | Owns the W0a config surface of the UVA weight-offload mirror floor per `.agents/specs/weight-offload-uva.md`. Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | +| `CLAIM-WEIGHT-OFFLOAD-W0A` | `ENG-WEIGHT-OFFLOAD` (`ACTIVE`) | see the matrix row and its spec — recorded here because the claim existed only as the row's owner cell and had no claims-directory file (record-rot repair, row/RECORD-ROT-REPAIR) | per the owning row's record | `row/ENG-WEIGHT-OFFLOAD` | Owns the W0a config surface of the UVA weight-offload mirror floor per `.agents/specs/weight-offload-uva.md`. Scope, evidence and exclusions are exactly the owning matrix row and its spec; this file only restores the claims-directory presence the gate requires. | `ACTIVE` | 2026-09-23 — claim file created by the record-rot repair; no scope change | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 152b527c11..835adf7876 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -56,15 +56,15 @@ 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 FINE-GRAINED half of W3 (a `hash_block_size` STRICTLY FINER than a group's block size, which still needs the `KV-BLOCK-POOL` align path that throws at `block_pool.cpp:93,220`) together with the mamba partial-tail stop, and the benchmark (W4) remain deferred. **The ENGINE-SIDE half of W3 LANDED 2026-09-12**, as the fix for a `vllm serve` abort on DeepSeek-V4 (`ISSUE-LOCAL-01M2BXYS2R8ZNC68QYGHK4RVQH`, row `MODEL-MM-deepseek-v4-deepseek-v4-for-causal-lm`): `LoadedEngine::ResolveSchedulerBlockSizes` re-derives the cache block size as the minimum group block size and calls this resolver, and the resulting pair is threaded through `MakeScheduler` into `Scheduler` / `AsyncScheduler`, mirroring `vllm/v1/engine/core.py:335-338` then `:158-170` @ `e126687a9a`. **The scheduler therefore NO LONGER hardcodes `hash_block_size = block_size`.** The default dense path stays byte-identical: one group whose block size equals the configured one resolves to the same pair it was passed before. | 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-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 `src/vllm/entrypoints/openai/server_main.cpp:622`; **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 FINE-GRAINED half of W3 (a `hash_block_size` STRICTLY FINER than a group's block size, which still needs the `KV-BLOCK-POOL` align path that throws at `block_pool.cpp:93,220`) together with the mamba partial-tail stop, and the benchmark (W4) remain deferred. **The ENGINE-SIDE half of W3 LANDED 2026-09-12**, as the fix for a `vllm serve` abort on DeepSeek-V4 (`ISSUE-LOCAL-01M2BXYS2R8ZNC68QYGHK4RVQH`, row `MODEL-MM-deepseek-v4-deepseek-v4-for-causal-lm`): `LoadedEngine::ResolveSchedulerBlockSizes` re-derives the cache block size as the minimum group block size and calls this resolver, and the resulting pair is threaded through `MakeScheduler` into `Scheduler` / `AsyncScheduler`, mirroring `vllm/v1/engine/core.py:335-338` then `:158-170` @ `e126687a9a`. **The scheduler therefore NO LONGER hardcodes `hash_block_size = block_size`.** The default dense path stays byte-identical: one group whose block size equals the configured one resolves to the same pair it was passed before. | 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` | - | @@ -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 `"