diff --git a/.agents/coordination.md b/.agents/coordination.md
index badf7d50..96f6f388 100644
--- a/.agents/coordination.md
+++ b/.agents/coordination.md
@@ -118,6 +118,23 @@ without the selected contention proof for their entire run are discarded.
 
 ## Active claims
 
+**Canonical DONE-owner reachability repair (`KV-PREFIX-CACHE`,
+`SAMPLE-LOGPROBS`, `SPEC-DFLASH`,
+`MODEL-SPEC-qwen3-dflash-dflash-qwen3-for-causal-lm`,
+`KERNEL-ATTN-DFLASH-BLOCK`, `KERNEL-ATTN-DFLASH-PAGED-BLOCK`, 2026-07-31,
+`CLAIM-RECORD-DONE-OWNER-REPAIR`).** Codex (GPT-5), isolated worktree
+`/home/mudler/.cache/sdd/localai-org-maint-bot-vllm.cpp/codex-cpu-regression-fix`,
+branch `codex/cpu-regression-fix`, base `upstream/main` `0d014016`. CPU-only
+governance repair: `scripts/check-agent-record.py` currently rejects six `DONE`
+rows because their three owner SHAs are not Git objects; replace them with the
+reachable commits that actually introduced the binding `DONE` closures. Owns
+only those six Owner cells, this claim, the minimal append-only state entry,
+and the required no-lifecycle-change checkpoint notes in `docs/STATUS.md` and
+`docs/BENCHMARKS.md`. No source, test, build, lifecycle, benchmark, README,
+roadmap, ledger, or support-claim change; verification is the agent-record
+checker and its mutation suite plus the remaining record/document checkers.
+No GPU, model, or compiler is required.
+
 **Stable C ABI contract spike (`SERVE-C-ABI`, 2026-07-31,
 `CLAIM-SERVE-C-ABI-SPIKE`).** Codex (GPT-5), isolated worktree
 `.worktrees/serve-c-abi-spike`, branch `codex/serve-c-abi-spike`, base
diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md
index 7a865327..9d060040 100644
--- a/.agents/engine-matrix.md
+++ b/.agents/engine-matrix.md
@@ -54,7 +54,7 @@ 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) | `444ea9d7` |
+| `KV-PREFIX-CACHE` | APC hashes, lookup, allocation, partial blocks, eviction, plus explicit/model-default cache policy. W0 ports arbitrary-group no-prefix coordination and makes hybrid/attention-free defaults cache-off. **Full-surface re-audit 2026-07-22 ([spike](specs/prefix-prompt-caching-parity.md)) — the ported core is DEEPER than this row read (chain hashing, pool, all three coordinators, the complete hybrid intersection and four single-type managers), and the residual gaps are narrower and DIFFERENT:** **`generate_block_hash_extra_keys`: W2 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC`)** — the hardcoded no-op is replaced by a 1:1 port of `kv_cache_utils.py:451-591` (`_gen_mm_extra_hash_keys` + LoRA name + `cache_salt`, fixed order lora->mm->salt; prompt_embeds deferred, no prompt-embeds path). `Request`/`EngineCoreRequest` carry `cache_salt` + `lora_name`; `FromEngineCoreRequest` sets them BEFORE the first hash (fixed a latent ordering bug: mm_features were assigned after the ctor already hashed). The latent correctness trap is CLOSED and RED-first proven: with the stub, a tenant-B request false-hits tenant-A's 48 cached tokens (`n1==48`); with extra keys `n1==0` (no false-share). This unblocks the MM + LoRA cache consumers. **prefix-cache statistics: CLOSED 2026-07-22** (W1) — `PrefixCacheStats`/`CachingMetrics` ported 1:1 with `log_stats` DEFAULTED ON, which unblocks the `BACKEND-GATE-CUDA-SGLANG-PREFIX` hit-proof requirement; first measured hit rate 0.75 on a repeated-prefix corpus; no `cache_salt`; 1 of upstream's 4 hash algos; `skip_reading_prefix_cache` absent; partial-block primitives throw (upstream's own are DEAD CODE — no caller in `vllm/` — so they are NOT owed as live behaviour). **Also cleared: the "blocked on a supported non-hybrid family" blocker is STALE** — dense models default APC ON and five have landed, yet NO gate has ever run cache-ON **MLA prefix-cache-hit assert fixed 2026-07-23** (`CLAIM-MLA-PREFIX-CACHE-ASSERT`): `FullAttentionManager::find_longest_cache_hit` asserted `kind()==kFullAttention`, aborting DeepSeek-V2 (MLA group, kind `kMlaAttention`, APC default-ON) under asserts-enabled builds — latent since `ec6f4be`, inert under Release/NDEBUG. Relaxed to upstream's precondition `isinstance(spec, FullAttentionSpec or ChunkedLocalAttentionSpec)` (single_type_kv_cache_manager.py:578-582; MLAAttentionSpec IS-A FullAttentionSpec) ⇒ accept `kFullAttention` / `kMlaAttention` / `kChunkedLocalAttention`; restores DeepSeek-V2 SACRED gate 8/8 asserts-on, full-attention byte-identical, new MLA prefix-cache-hit unit cases. | T0 | `vllm/config/model.py:1805-1860`; `vllm/engine/arg_utils.py:510,1160-1166,2473-2508`; `vllm/config/cache.py:39,93,95`; extra keys `vllm/v1/core/kv_cache_utils.py:539-574`; hasher factory `:673-730`; `vllm/v1/core/kv_cache_coordinator.py:377-425,782-834`; `tests/v1/core/test_prefix_caching.py:225,1475,2781` | hashes/managers `src/vllm/v1/core/kv_cache_utils.cpp:259,291`; **extra_keys** `generate_block_hash_extra_keys` + `_gen_mm_extra_hash_keys` `src/vllm/v1/core/kv_cache_utils.cpp`; `cache_salt`/`lora_name` on `include/vllm/v1/request.h` + `include/vllm/v1/engine/types.h`, copied in `src/vllm/v1/request.cpp` `FromEngineCoreRequest` (fields set before the first hash); `src/vllm/v1/core/kv_cache_manager.cpp:124`; no-prefix coordinator/factory `src/vllm/v1/core/kv_cache_coordinator.cpp:260,273,279,545`; model-default/hasher selection `src/vllm/entrypoints/model_loader.cpp:109,167,180,191`; CLI `examples/server/main.cpp:126`; **statistics** `include/vllm/v1/metrics/stats.h`, recorded `src/vllm/v1/core/kv_cache_manager.cpp:139-147`, reset flag `:270-276`, take-and-swap `make_prefix_cache_stats()`, per-step window fold at the end of `Scheduler::schedule()`, accessors `Scheduler`/`EngineCore`/`LLMEngine::prefix_cache_metrics()`; `Request::num_preemptions` un-deferred (`include/vllm/v1/request.h`, incremented in `Scheduler::preempt_request`) | existing APC primitives `tests/vllm/v1/test_kv_cache_utils.cpp:411,516,536`; no-prefix hybrid allocation/no-hit `tests/vllm/v1/test_kv_cache_coordinator.cpp:213`; default/override resolution `tests/vllm/entrypoints/test_loaded_engine_dense.cpp:343`; server help and online cache-off contracts `examples/CMakeLists.txt:34`; `tests/tools/test_online_gate_client.py:582,633`; statistics plus the first MEASURED hit rate `tests/vllm/v1/test_prefix_cache_stats.cpp` 12/12; **W2 extra_keys** — ported mm/lora/salt cases + ordering + hash-level no-false-share `tests/vllm/v1/test_kv_cache_utils.cpp` (29/29), manager-level salt-partition no-false-share (RED-proven `n1 48->0`) `tests/vllm/v1/test_kv_cache_manager.cpp` (10/10), CPU gate on dgx GB10. **W3 DONE 2026-07-27 (`CLAIM-ROADMAP-D4APC-W3`, dgx GB10, NOT pushed) — the FIRST-EVER cache-ON model gate:** `tests/parity/test_qwen3_apc_e2e.cpp` on `Qwen/Qwen3-4B` (dense, full-attention, APC-default-ON) 2/2 cases, 84/84 asserts — APC-ON hits 2240/2777 (rate 0.807) / APC-OFF 0; APC-ON == APC-OFF token-exact 5/6 (1 diff a vLLM-confirmed 0.125-nat near-tie); == vLLM-APC-ON teacher-forced (OFF 6/6 gap 0.0, ON 6/6 gap ≤0.125 nats, 0 outside top-20); TTFT 70.1→39.9 ms = 1.76×. NO engine code changed (gate-only over the already-shipped default-ON path); 4B SACRED 16/16 no-regression. Oracle vLLM 0.25.0. Ledger: [parity-ledger.md#L746](parity-ledger.md#L746) | [prefix-prompt-caching-parity.md](specs/prefix-prompt-caching-parity.md) (umbrella); [prefix-caching.md](specs/prefix-caching.md) (cache-policy leaf) | `DONE` (dense APC path; W4 events/W5 partial/W6 mamba-align/W7 reset endpoint tracked in `KV-EVENTS`/`KV-MAMBA-ALIGN`/own future rows) | `a41af480` |
 | `KV-PREFIX-MATCH-UNIT` | `--prefix-match-unit` (config `prefix_match_unit`): the finest token boundary a prefix-cache hit can land on == the `hash_block_size`/"prefix match unit" the block hasher uses. NEW in 0.26 (absent at the prior `e24d1b24`/0.25.0 pin). For a HYBRID/multi-group model the resolver `resolve_kv_cache_block_sizes` computes `hash_block_size = prefix_match_unit if set else gcd(group_block_sizes)` (scheduler block size = `lcm`), letting matching land FINER than a physical block (e.g. 16/32 tokens inside a 1024-token block) provided every group block size is divisible by it; single-group (dense) models ignore the knob. Backs off to the scheduler block size when no prefix-cache/connector consumer is active or a mamba group diverges from `cache_block_size` (mamba_cache_mode != "align"); throws on a non-divisible unit. **W0 spike + W1 resolver LANDED 2026-07-28 (`CLAIM-PREFIX-MATCH-UNIT`, NOT pushed):** `resolve_kv_cache_block_sizes` ported 1:1 (explicit-parameter signature vs upstream's `VllmConfig`, our config surface is threaded), RED-first unit-gated (default gcd `!=` `=16` override). `PARTIAL`: the config/CLI/ABI field (W2), the scheduler threading of a resolved `hash_block_size != block_size` + mamba partial-tail stop (W3, needs the `KV-BLOCK-POOL` align path that still throws), and the benchmark (W4) are deferred. Default path byte-identical (single-group inert; scheduler still passes `block_size`). | T1 | `vllm/engine/arg_utils.py:696,1222,1940`; `vllm/config/cache.py:56-67`; resolver `vllm/v1/core/kv_cache_utils.py:626-688`; hasher `:691-748`; call site `vllm/v1/engine/core.py:154`; scheduler `vllm/v1/core/sched/scheduler.py:76,268-270,282,312-318`; fine-grained view `vllm/v1/core/single_type_kv_cache_manager.py:683,697` | resolver `src/vllm/v1/core/kv_cache_utils.cpp:638` (`resolve_kv_cache_block_sizes`), decl `include/vllm/v1/core/kv_cache_utils.h`; hash_block_size already plumbed `get_request_block_hasher` `src/vllm/v1/core/kv_cache_utils.cpp:577`; DEFERRED align path throws `src/vllm/v1/core/block_pool.cpp:93,220` (shared with `KV-BLOCK-POOL`) | `tests/vllm/v1/test_prefix_match_unit.cpp:64,88,99,119,129,145,164,186` 8/8 (29 assertions): single-group inert + DCP scale, multi-group default=gcd, `=16` override finer-than-default (RED), finer-than-1024-block, non-divisible throws, no-consumer back-off + connector-alone re-enable, mamba non-align back-off vs align gcd, hasher-granularity RED (coarse 2 vs fine 4 hashes); [parity-ledger.md](parity-ledger.md) | [prefix-match-unit.md](specs/prefix-match-unit.md) | `PARTIAL` | `CLAIM-PREFIX-MATCH-UNIT` |
 | `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` | - |
@@ -128,7 +128,7 @@ lifecycle are unchanged.
 |---|---|---|---|---|---|---|---|---|
 | `SAMPLE-CORE` | Ordered temperature, top-k/p, min-p, penalties, seed, stop, length, output-kind pipeline | T0 | `vllm/v1/sample/sampler.py:20,72,243`; `vllm/sampling_params.py:264,500`; `vllm/v1/worker/gpu_input_batch.py:889-963`; `tests/v1/sample/test_sampling_params_e2e.py:17,25,40,176` | `src/vllm/sampling_params.cpp:25,35,167`; `src/vllm/v1/sample/sampler.cpp:152,215`; `src/vllm/v1/worker/gpu/input_batch.cpp:255,344,451`; `src/vllm/v1/core/sched/utils.cpp:12` | `tests/vllm/test_sampling_params.cpp:13,64,231,265`; `tests/vllm/v1/sample/test_sampler.cpp:46,78,118,142,165,253`; `tests/vllm/v1/worker/test_input_batch.cpp` (C7 wiring: min_p/min_tokens/logprobs-count reach SamplingMetadata + condense/swap); `tests/vllm/v1/test_input_processor.cpp` (all_stop_token_ids) | [sampling-controls-c7.md](specs/sampling-controls-c7.md) (`SAMPLE-CORE`) | `ACTIVE` | `CLAIM-ROADMAP-C7` |
 | `SAMPLE-PHILOX` | Torch-Philox bit-exact stochastic parity | T1 | `vllm/v1/sample/ops/topk_topp_sampler.py:70`; `vllm/v1/sample/sampler.py:243` | - | - | `planned: specs/philox-rng-parity.md` | `INVENTORIED` | - |
-| `SAMPLE-LOGPROBS` | Token logprobs payload end to end | T1 | `vllm/logprobs.py:12,157,175`; `vllm/v1/engine/logprobs.py:29,69,348`; `vllm/v1/outputs.py:28,38`; `vllm/v1/core/sched/scheduler.py:1815-1836`; `vllm/entrypoints/openai/completion/serving.py:652`; `vllm/entrypoints/openai/chat_completion/serving.py:1114,1141`; `tests/v1/sample/test_logprobs.py:303` | `include/vllm/logprobs.h` (Logprob/LogprobsOnePosition/SampleLogprobs + AppendLogprobsForNextPosition); `include/vllm/v1/engine/logprobs.h`+`src/vllm/v1/engine/logprobs.cpp` (LogprobsProcessor); `src/vllm/v1/outputs.cpp:24` (LogprobsTensors::slice_request); `src/vllm/v1/worker/gpu/runner.cpp:1272` (ModelRunnerOutput.logprobs); `src/vllm/v1/core/sched/scheduler.cpp:674` (slice -> EngineCoreOutput.new_logprobs); `src/vllm/v1/engine/output_processor.cpp` (LogprobsProcessor integration -> CompletionOutput.logprobs); `src/vllm/entrypoints/openai/serving_utils.cpp` (BuildCompletionLogProbs/BuildChatLogprobs); `src/vllm/entrypoints/openai/protocol.cpp` (CompletionLogProbs/ChatCompletionLogProbs to_json); `src/vllm/entrypoints/openai/serving_{completion,chat}.cpp` | `tests/vllm/entrypoints/openai/test_logprobs.cpp:67,103` (serialization vs vLLM oracle, RED-first N vs N+1; LogprobsProcessor accumulation + inertness `:137,167`); `tests/vllm/entrypoints/openai/test_serving.cpp:411,668` (e2e through the CPU engine: logprobs=K + chat top_logprobs, inert-when-off `:467`). Closing record [parity-ledger.md#L741](parity-ledger.md#L741) 2026-07-27 W5 row | [sampling-controls-c7.md](specs/sampling-controls-c7.md) (W5) | `DONE` | `7a3f04b2` |
+| `SAMPLE-LOGPROBS` | Token logprobs payload end to end | T1 | `vllm/logprobs.py:12,157,175`; `vllm/v1/engine/logprobs.py:29,69,348`; `vllm/v1/outputs.py:28,38`; `vllm/v1/core/sched/scheduler.py:1815-1836`; `vllm/entrypoints/openai/completion/serving.py:652`; `vllm/entrypoints/openai/chat_completion/serving.py:1114,1141`; `tests/v1/sample/test_logprobs.py:303` | `include/vllm/logprobs.h` (Logprob/LogprobsOnePosition/SampleLogprobs + AppendLogprobsForNextPosition); `include/vllm/v1/engine/logprobs.h`+`src/vllm/v1/engine/logprobs.cpp` (LogprobsProcessor); `src/vllm/v1/outputs.cpp:24` (LogprobsTensors::slice_request); `src/vllm/v1/worker/gpu/runner.cpp:1272` (ModelRunnerOutput.logprobs); `src/vllm/v1/core/sched/scheduler.cpp:674` (slice -> EngineCoreOutput.new_logprobs); `src/vllm/v1/engine/output_processor.cpp` (LogprobsProcessor integration -> CompletionOutput.logprobs); `src/vllm/entrypoints/openai/serving_utils.cpp` (BuildCompletionLogProbs/BuildChatLogprobs); `src/vllm/entrypoints/openai/protocol.cpp` (CompletionLogProbs/ChatCompletionLogProbs to_json); `src/vllm/entrypoints/openai/serving_{completion,chat}.cpp` | `tests/vllm/entrypoints/openai/test_logprobs.cpp:67,103` (serialization vs vLLM oracle, RED-first N vs N+1; LogprobsProcessor accumulation + inertness `:137,167`); `tests/vllm/entrypoints/openai/test_serving.cpp:411,668` (e2e through the CPU engine: logprobs=K + chat top_logprobs, inert-when-off `:467`). Closing record [parity-ledger.md#L741](parity-ledger.md#L741) 2026-07-27 W5 row | [sampling-controls-c7.md](specs/sampling-controls-c7.md) (W5) | `DONE` | `ecda3ce1` |
 | `SAMPLE-PROMPT-LOGPROBS` | Prompt logprobs (payload plumbing + serialization done; runner prompt-position logits SOURCE pending) | T1 | `vllm/v1/engine/logprobs.py:121`; `vllm/v1/worker/gpu_model_runner.py:3842` (`_get_prompt_logprobs_dict`); `vllm/sampling_params.py:303` | payload path landed: `src/vllm/v1/engine/logprobs.cpp:75,100` (`UpdatePromptLogprobs`/`pop_prompt_logprobs`); `include/vllm/v1/engine/types.h:129,161` (`ModelRunnerOutput.prompt_logprobs_dict`, `EngineCoreOutput.new_prompt_logprobs_tensors`); `src/vllm/v1/core/sched/scheduler.cpp:688` (prompt_logprobs_dict slice); `src/vllm/v1/engine/output_processor.cpp:224` (`RequestOutput.prompt_logprobs`). **RESIDUAL: the runner does not yet compute prompt-position logits (lm_head over prompt tokens) — the tensor SOURCE; a runner/prefill addition adjacent to C5** | `tests/vllm/entrypoints/openai/test_logprobs.cpp:137,167` (LogprobsProcessor accumulation/inertness cover the shared consume path; the prompt-only source is inert until the runner lands) | `planned: specs/prompt-logprobs.md` | `PARTIAL` | - |
 | `SAMPLE-LOGPROB-TOKEN-IDS` | `logprob_token_ids` generative scoring and `logprobs_mode` variants beyond raw/processed (`SAMPLE-LOGPROBS` covers the payload only); carried from porting-inventory §6 (T1) at the v1 fold | T1 | `vllm/sampling_params.py:278,727`; `vllm/config/model.py:82,221` | - | - | `planned: specs/logprob-token-ids.md` | `INVENTORIED` | - |
 | `SAMPLE-LOGIT-FILTERS` | Logit bias, allowed-token IDs, bad words | T1 | `vllm/sampling_params.py:318,321,337,341,388-413,659-698`; `vllm/v1/sample/sampler.py:396`; `vllm/v1/worker/gpu_input_batch.py:446-471`; `vllm/entrypoints/openai/completion/protocol.py:369-371`; `tests/v1/sample/test_sampler.py:367,413`; `tests/v1/sample/test_sampling_params_e2e.py:106,147` | `include/vllm/sampling_params.h`+`src/vllm/sampling_params.cpp` (fields+validation); `src/vllm/entrypoints/openai/protocol.cpp` (`ParseLogitFilters`/`ApplyLogitFilters` clamp); `src/vllm/v1/worker/gpu/input_batch.cpp:255,344` (per-slot wiring+condense/swap); `src/vllm/v1/engine/input_processor.cpp` (bad_words tokenization); `src/vllm/v1/sample/sampler.cpp:239`; `src/vllm/v1/sample/logits_processor/builtin.cpp:41`; `src/vllm/v1/sample/ops/bad_words.cpp:13,55` | `tests/vllm/v1/sample/test_logits_processors.cpp:121,163,200`; `tests/vllm/test_sampling_params.cpp` (bad_words/allowed_token_ids validation); `tests/vllm/entrypoints/openai/test_protocol.cpp` (logit_bias clamp+parse); `tests/vllm/v1/worker/test_input_batch.cpp` (wiring, RED-first); `tests/vllm/v1/test_input_processor.cpp` (bad_words tokenization) | [sampling-controls-c7.md](specs/sampling-controls-c7.md) (`SAMPLE-LOGIT-FILTERS`) | `ACTIVE` | `CLAIM-ROADMAP-C7` |
@@ -163,7 +163,7 @@ lifecycle are unchanged.
 | `SPEC-DFLASH-GGUF` | DFlash speculative decoding from GGUF, two axes: (A) GGUF DRAFT + safetensors target, (B) GGUF target too. llama.cpp master carries a full `dflash` GGUF contract (arch string `dflash`, tensors `fc`/`enc.output_norm`/`output_norm`/`blk.N.*`, KVs `dflash.target_layers` + `dflash.target_hidden_size`); the arch is ABSENT from checkouts older than ~2026-07, so a stale tree reads as "no contract exists". The GGUF tensor set omits `token_embd`/`output` because the draft SHARES the target's embed+lm_head, which is exactly what `LoadDflashDraft` already does. Blockers are in the loader, not the model: `MakeDflashDraftConfig` reads `draft_dir/config.json` (a GGUF has none), `ResolveDflashDraftDir` probes for `config.json` so it cannot see a `.gguf`, and `LoadDflashDraft` is typed on `std::vector<SafetensorsFile>` for the shared bf16 head (the axis-B blocker). Axis A independently shippable. NO ABI change | T2 | llama.cpp `origin/master` @ 2026-07-28 (tag era `b10158`): `gguf-py/gguf/constants.py:547,1151,4350`; `gguf-py/gguf/tensor_mapping.py:1297-1305` (`ENC_OUTPUT_NORM`<-`model.hidden_norm`, `FC`<-`model.fc`); `conversion/qwen.py:351` (mask token via the standard tokenizer KV); `convert_hf_to_gguf.py --target-model-dir` | **GD1-GD7 LANDED 2026-07-28 (BOTH AXES COMPLETE and PROVEN end to end on GB10)**: `MakeDflashGgufConfig` + `LoadQwen3DFlashFromGguf` `src/vllm/model_executor/models/qwen3_dflash_gguf.cpp:88,227` (+ header), `IsDflashGgufDraft` + the `.gguf` branch in `ResolveDflashDraftDir`/`LoadDflashDraft` `src/vllm/entrypoints/model_loader.cpp:121,222`. Goes through the `TensorResolver` seam (unlike `SPEC-MTP-GGUF`) because dflash norms are RAW, so the existing `LoadQwen3DFlash` qkv/gate_up concatenation is reused unchanged. **`GD4` defect FIXED** (`model_loader.cpp:238-249`): the GGUF branch left `config.vocab_size` 0 - correct for `MakeDflashGgufConfig` (the DFLASH arch has no vocab KV and no `token_embd`) but fatal for the forward, which sizes the shared embedding view as `{config.vocab_size, H}`, so the first propose threw `cuda embedding: empty table (vocab 0)`. Now back-filled from the target's `embed_tokens` rows (the condition is on the VALUE, not the draft source, so it generalizes to a GGUF target). Load-level green had hidden it; only GENERATING found it. **GD5-GD7 = axis B**: `SharedHeadSource` `src/vllm/entrypoints/model_loader.cpp` re-expresses the shared bf16 `embed_tokens`+`lm_head` seam as a SOURCE and re-types `LoadDflashDraft`'s second parameter - THAT TYPE was the whole axis-B blocker - with the GGUF arm `LoadGgufSharedEmbedAndHeadBf16` `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:773` reusing the trunk loader's tied-embedding rule and sidecar-aware dequant instead of restating them; the shared-head load moved into ONE common tail so all four (draft format x target container) combinations run identical code; the `dflash` half of the GGUF-branch rejection `model_loader.cpp` is deleted (the `mtp` half untouched) and the draft load is wired into the GGUF branch | `tests/vllm/models/test_qwen3_dflash_gguf.cpp:36,84` 2 cases / 47 assertions against the REAL published Qwen3.6-27B DFlash draft (env-gated `VLLM_DFLASH_GGUF_MODEL`, CI asset-free): the +1 target-layer offset undone against the KV read back from the same file, block_size/mask_token present, vocab_size left 0, layer_types cover every block, fc `[H, H*num_taps]` with `nk` SET, qkv/gate_up row-concat shapes, embed/lm_head left EMPTY for the target. **RED-first BEHAVIOURAL** (dropping the `-1` fails the offset checks). **`GD4` e2e gate** `tests/parity/test_qwen27_dflash_spec_decode.cpp:343` (second case, draft source env-driven via `VLLM_DFLASH_DRAFT`/`_B`; asset-gated, CI-inert): on dgx GB10 sm_121a against the Qwen3.6-27B NVFP4 safetensors target, the Q4_K_M GGUF draft and the bf16 z-lab safetensors draft produce **token-for-token IDENTICAL** DFlash-ON continuations with **IDENTICAL** accepted/proposed (20/80 on a 24-token prompt, 42/96 on a 48-token prompt), spec-OFF self-reproducible 3/3 and 0 exact ties (min margin 0.197/0.400 nats). Regression: gguf_mtp 19, qwen35_gguf_spec_decode 10, gguf 103, gguf_qwen36_loader 99, gguf_keep_quant 5958, ops_gdn 1825, llm_engine 196, capi 232, runner 257 all unchanged. **`GD5` unit gate** `tests/vllm/test_gguf_qwen36_loader.cpp` 3 new synthetic-GGUF cases (6 cases / 286 assertions total, CPU and the dgx CUDA build): the untied head really comes from `output.weight` and not the embedding (distinct fill values), the tied fallback aliases it onto `token_embd`, the `nk` flags separate the gather table from the MatmulBT weight, a file with no `token_embd` is refused. 3-mutant battery, 3 caught (`nk` flipped, head forced to the embedding, tied forced false). **`GD7` e2e gate** `tests/parity/test_qwen27_dflash_spec_decode.cpp` third case (targets env-driven via `VLLM_DFLASH_TARGET_B`; asset-gated, CI-inert): on dgx GB10 sm_121a the Qwen3.6-27B NVFP4 **GGUF** target + `Q4_K_M` GGUF draft loads, takes the shared head from the GGUF, generates, and its DFlash-ON continuation is **token-for-token IDENTICAL to that same target's spec-OFF** (24/24, the STRICT form) with acceptance ALIVE at 14/160; 1 case / 15 assertions, exit 0. **The spike's highest risk is EMPTY on this asset, proven not assumed**: the 27B NVFP4 GGUF stores `token_embd`/`output` as ggml BF16, byte-identical to the safetensors sibling (2,542,796,800 bytes each, ZERO differing), so B1's shared-head read is verbatim, not a dequant. Acceptance IS lower than the safetensors-target arm and is NOT chargeable to the head: the two containers diverge at index 4 with NO speculation, because `QUANT-GGUF-NVFP4` is dequant-only so the GGUF target computes in bf16 while the safetensors target runs the true W4A4 kernels. **RE-MEASURED 2026-07-29 on a PRODUCTION-CONFIGURED build (`CLAIM-GGUF-SPEC-REVERIFY`), because every GD4/GD7 GPU number above came from a build configured WITHOUT `-DVLLM_CPP_CUTLASS_DIR` and WITHOUT `-DVLLM_CPP_TRITON=ON`.** Build proven correct three ways (see the `SPEC-MTP-GGUF` row; SACRED 27B **235/235**, `cuobjdump` 40 cubins all `sm_121a`). **AXIS B HOLDS EXACTLY**: `test_qwen27_dflash_spec_decode -tc="dflash axis-B*"` **15/15 assertions, exit 0**, GGUF-target DFlash-ON token-identical to that target's own spec-OFF 24/24, acceptance **14/160 unchanged**, cross-target spec-OFF divergence still at index 4, 81.01 GiB peak RSS, 6m53.08s. **AXIS A WAS RED ON THE 48-TOKEN PROMPT (reproducibly, 3 of 3 runs) AND IS NOW CLOSED.** The RED was real: cross-format TOKEN identity held on both prompts, but the exact accept-count half of bar (a) failed (`arm_a.proposed == arm_b.proposed` / `arm_a.accepted == arm_b.accepted`) because the Q4_K_M draft measured **46/112** against the bf16 z-lab draft's **47/96** (one extra 16-wide propose block, one fewer acceptance, zero token difference), 15/17, exit 1; the 24-token prompt stayed green at 17/17 with both drafts at 15/144. **`GD9` 2026-07-29 root-caused it IN WEIGHT SPACE as ordinary `Q4_K_M` cost, category (a), not a defect in our GGUF draft path - and the bar's own premise ("Same weights, two containers") was false for the asset it was pointed at.** The publishing repo also carries an UNQUANTIZED `BF16` GGUF (3,471,497,440 B) beside `Q8_0`/`Q6_K`/`Q5_K`/`Q4_K_M`, which the spec had recorded as nonexistent; that retired the `NOT APPLICABLE` on gate 2. CPU gate `tests/vllm/models/test_qwen3_dflash_gguf.cpp` third case (asset-gated `VLLM_DFLASH_GGUF_BF16_MODEL` + `VLLM_DFLASH_ST_DIR`): `LoadQwen3DFlashFromGguf(BF16)` is **BYTE-IDENTICAL to `LoadQwen3DFlash(z-lab shards)` on all 58 tensors, 302/302 assertions, exit 0**, and FUNCTIONALLY RED against the `Q4_K_M` file (21/302 red, exactly the 21 quantized matmul tensors), so not a vacuous pass. Supporting: our `DequantGgufRowToBf16` is bit-equal to `gguf-py`'s `gguf.quants.dequantize` on the real `fc.weight` (Q4_K), `blk.0.attn_q.weight` (Q4_K) and `blk.2.ffn_down.weight` (Q6_K), zero differing bf16 values; the ladder's mean relative weight error is monotone and uniform with NO outlier tensor (BF16 0, Q8_0 5.6e-3, Q6_K 1.85e-2, Q5_K 3.85e-2, Q4_K_M 7.6e-2); the only numeric config delta is `rms_norm_eps` at 2.5e-9 relative. Also landed: an off-by-default `VT_SPEC_TRACE=1` per-block propose/accept trace in `GPUModelRunner::sample_tokens_with_rejection` (`src/vllm/v1/worker/gpu/runner.cpp`). **`GD10` 2026-07-29 CONFIRMED IT END TO END ON GB10 and closed gates 3 and 5.** Build proven production-configured three ways (configure log 0 `CUTLASS not found`; `cuobjdump -lelf` 40 cubins ALL `sm_121a` zero `sm_75` on both binaries; SACRED `test_qwen27_paged_engine` **235/235, exit 0**, 31.34s, 23.67 GiB). The **`BF16` GGUF draft reads EXACTLY 47/96**, the safetensors draft's own number, at 48 tokens on the discriminating prompt - reproduced 2 of 2 - plus 27/64 = 27/64 at 24 tokens and 15/144 = 15/144 on the second prompt, tokens IDENTICAL throughout, 17/17 exit 0 each time; the `Q4_K_M` arm reads 46/112 on the SAME binary in the SAME `flock` series. Restoring only the draft's numeric precision restores the count, so quantization is the whole cause and nothing structural survives. Bar (a) is consequently SPLIT rather than relaxed (`tests/parity/test_qwen27_dflash_spec_decode.cpp`): tokens stay EXACT unconditionally; accept counts are EXACT on a cross-FORMAT arm and BANDED (`abs(d_accepted) <= 2`, `abs(d_proposed) <= k*2`) on a cross-QUANTIZATION one, with the arm chosen by `IsQuantizedGgufDraft` reading the draft file's ggml types (`GgmlTraits().block_elems > 1`) rather than by a flag. The band is derived, not picked: measured `d_accepted` is 0, 0, -1, so the bound is that maximum plus one quantum; and `d_proposed = -k * d_accepted` EXACTLY once the token streams match (confirmed at -1 / +16), so the proposed bound follows. **Mutation-proved non-vacuous**: rebuilt at band 0 the `Q4_K_M` arm is 15/17 exit 1 while the `BF16` arm stays 17/17 exit 0 on the exact branch. **AXIS B BROADENED from ONE prompt to THREE**, strict form green on all: "The capital of France is" IDENTICAL 14/160 (15/15), "Write a Python function that reverses a string:" IDENTICAL 24/64 (15/15), "Photosynthesis is the process by which" IDENTICAL 15/128 (9/9), all exit 0, ~6m30-6m52 and ~81 GiB peak RSS each. The second prompt REFINES the recorded acceptance claim: the safetensors-target arm is ALSO 24/64 there with the two containers' DFlash-ON streams IDENTICAL, so the GGUF target's lower acceptance is prompt-dependent (their spec-OFF streams diverge at index 4 on the first prompt, index 16 on the second) and not a standing penalty; the cause remains `QUANT-GGUF-NVFP4` being dequant-only, with the shared head excluded by a byte comparison. Gates 1-5 and 7 MET; gate 6 (speed) `PENDING` BY DESIGN and not owed - a DFlash-ON throughput A/B between the two target containers is not a fair comparison until a native NVFP4 GGUF GEMM exists. Evidence: [docs/BENCHMARKS.md](../docs/BENCHMARKS.md) top section, [parity-ledger.md](parity-ledger.md#L845) | [specs/gguf-dflash-draft.md](specs/gguf-dflash-draft.md) | `DONE` | `c62f2fa3` |
 | `SPEC-REJECTION` | Rejection sampler. **I3 verify half LANDED (2026-07-24)**: per-request logits EXPANSION to `1 + k_i` rows (`StepInputs::cu_num_logits` / `num_draft_tokens_per_req` / expanded `logits_indices`) plus the GREEDY rejection sampler — accept a draft iff it equals the target argmax at its own position, emit the target argmax on the FIRST mismatch and stop, emit the bonus argmax when all `k_i` accept, `num_sampled = accepted + 1`, `num_rejected = k_i - accepted` (feeds I2's `num_computed_tokens` rollback and `InputBatch::num_accepted_tokens`). One additive vt op (`kGreedyRejectionSample`) with a CPU reference and a CUDA two-phase mirror of upstream's row-argmax + one-thread-per-request accept walk. DEFAULT-OFF and INERT: with no `SpeculativeConfig` no drafts are ever scheduled, `cu_num_logits` is `arange(num_reqs+1)`, `logits_indices` is the pre-change array and the runner never enters the rejection branch. STOCHASTIC/Gumbel, block verification, `apply_sampling_params` over the expanded batch, and the spec grammar bitmask stay DEFERRED (M-mtp-3). **I5b DRAFTER PREFILL INPUT-PREP LANDED (2026-07-24, `CLAIM-SPEC-MTP-I5B`)**: the draft-token input splice this row's I3 note deferred to I5 — `vllm::v1::prepare_prefill_inputs` + its `SpecPrefillInputs` output struct shift each request's `input_ids` left one within its query span, splice the just-sampled next token (`num_sampled>0 ? last_sampled[idx_mapping[r]] : next_prefill_tokens[...]`) into the freed slot, `query_len -= num_rejected`, and emit last-token index / query_start_loc / seq_lens + CG padding (mirror `speculator.py:469-588`, k=1 early-exit :236-238). A HOST routine in a NEW spec_decode-tree TU (no new CUDA kernel; mirrors the DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family — the DGX runner leaf ports the loop to the Triton kernel at I5d), unit-gated `test_prepare_prefill_inputs` 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT (nothing calls it until I5d), additive by construction. Row stays `ACTIVE` — the e2e greedy token gate (M-mtp-1) is owed before `DONE` | T1 | `vllm/v1/worker/gpu/spec_decode/rejection_sampler.py:43,101-160`; `rejection_sampler_utils.py:524,564-585,628,828-841,846-849,863-1125`; `vllm/v1/worker/gpu/model_runner.py:866-898,1065-1077`; `vllm/v1/worker/gpu/input_batch.py:303-397,408-453`; **I5b** `vllm/v1/worker/gpu/spec_decode/autoregressive/speculator.py:469-588,236-238` | `include/vllm/v1/spec_decode/rejection_sampler.h`; `src/vllm/v1/spec_decode/rejection_sampler.cpp`; `include/vt/ops.h` (`kGreedyRejectionSample`, `vt::GreedyRejectionSample`); `src/vt/cpu/cpu_sample.cpp` (CPU reference); `src/vt/cuda/cuda_sample.cu` (`RejectionRowArgmaxKernel` + `GreedyRejectAcceptKernel`); `src/vt/ops.cpp`; `include/vllm/v1/worker/gpu/prepare_inputs.h` + `src/vllm/v1/worker/gpu/prepare_inputs.cpp` (the expansion); `include/vllm/v1/worker/gpu/runner.h` + `src/vllm/v1/worker/gpu/runner.cpp` (`step_num_logits`, `sample_tokens_with_rejection`); **I5b** `include/vllm/v1/worker/gpu/spec_decode/autoregressive/prepare_prefill_inputs.h` + `src/vllm/v1/worker/gpu/spec_decode/autoregressive/prepare_prefill_inputs.cpp` | `tests/vllm/v1/spec_decode/test_rejection_sampler.cpp`; `tests/vllm/v1/worker/test_prepare_inputs.cpp` (expansion + no-draft byte-identity); `tests/vt/test_cuda_ops.cpp` (CUDA==CPU bit-exact at vocab 248320); **I5b** `tests/vllm/v1/spec_decode/test_prepare_prefill_inputs.cpp` (7 cases / 27 assertions, RED-first) | [mtp-spec-decode.md §2.4,§5](specs/mtp-spec-decode.md) | `ACTIVE` | `CLAIM-SPEC-REJECTION-I3`, `CLAIM-SPEC-MTP-I5B` |
 | `SPEC-GDN-SEGMENTS` | GDN speculative metadata and slot-snapshot rollback. **I4 LANDED (2026-07-24):** the spec/non-spec metadata split with decode→prefill reclassification (the #34845 case), the `T>1`/`IS_SPEC` GDN recurrence with per-timestep state snapshots, the conv sliding window advancing by the ACCEPTED count, and the k+1 state-slot allocation. DEFAULT-OFF and INERT (`num_spec==0` ⇒ `num_spec_decodes==0`, no shipped kernel branched — both spec kernels are NEW op ids). ROLLBACK PROVEN bit-exact: for every rejection point j the surviving SSM state and conv window are memcmp-identical to running only the accepted prefix through the shipped `vt::GdnDecode`/`CausalConv1dUpdate`, at the real 27B (Hv=48) and 35B (Hv=32) GDN dims on CPU and CUDA. MEASURED state cost: one f32 SSM slot = Hv·Dv·Dk·4B ⇒ 144 MiB/req (27B, 48 layers) / 60 MiB/req (35B, 30 layers) per extra slot; k=1 doubles the GDN SSM state. **I5a GDN LAYER ROUTING WIRED (2026-07-24, `CLAIM-SPEC-MTP-I5A`):** `GdnBlockPaged`'s `num_spec_decodes>0` branch now routes a PURE-spec batch through `vt::CausalConv1dSpecUpdate` + `vt::GdnSpecDecode` (mirror `qwen_gdn_linear_attn.py:1344-1357,1455-1475`), and the runner per-step upload (`StepDevInputs`/`BuildStepDevInputs` + the two decode-graph `Refresh` copies) now carries I4's six spec device tensors, gated by the extended `ValidateGdnAttentionMetadata` spec contract. DEFAULT-OFF INERT (`num_spec_decodes==0` ⇒ stub uploads + the identical non-spec branch). BIT-EXACT vs the I4 ops applied as a token-sequential decode chain, at the real 27B/35B GDN dims, via `GdnBlockPagedForTest` (`tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp`, CPU bit-exact + CUDA on-device); RED-first by a reverted stub (spec recurrence zeroed ⇒ 4/8 fail, maxΔ 1.3-1.6). MIXED spec+non-spec batch refused loudly — lands with I5d's runner loop. Row advances to `ACTIVE`: the M-mtp-1 e2e greedy token gate (verify/propose runner wiring) is owed before `DONE`, and `SPEC-MTP` STAYS `GATING` | T1 | `vllm/v1/attention/backends/gdn_attn.py:189-326,413-462`; `fla/ops/fused_sigmoid_gating.py:66-72,103-116,156-166`; `mamba/ops/causal_conv1d.py:818-1067,1181-1184`; `qwen_gdn_linear_attn.py:1329-1576`; `mamba_utils.py:213-234`; `mamba/abstract.py:55-59` | `include/vllm/v1/attention/backends/gdn_attn.h`; `src/vllm/v1/attention/backends/gdn_attn.cpp`; `include/vt/ops.h` (`kGdnSpecDecode`, `kCausalConv1dSpecUpdate`); `src/vt/ops.cpp`; `src/vt/cpu/cpu_ops.cpp`; `src/vt/cuda/cuda_gdn.cu`; `src/vllm/model_executor/models/qwen3_5_common.{h,cpp}` (`MakeQwen3_5KVCacheSpec`); **I5a:** `src/vllm/model_executor/models/qwen3_5.cpp` (`GdnBlockPaged` spec branch, `StepDevInputs`/`BuildStepDevInputs`, `ValidateGdnAttentionMetadata`), `src/vllm/model_executor/models/qwen3_5_internal.h` (`GdnBlockPagedForTest`) | `tests/vllm/v1/attention/test_gdn_metadata_builder.cpp` (20 cases / 483 assertions incl. the full upstream `GDN_BUILD_TEST_CASES` + default-off byte-identity); `tests/vt/test_ops_gdn.cpp` (reject-at-every-j rollback, CPU + CUDA, real dims); `tests/vllm/models/test_model_registry.cpp` (k+1 slot / widened-conv sizing + `num_spec==0` identity); **I5a** `tests/vllm/models/test_qwen3_5_gdn_spec_routing.cpp` (spec-routing bit-exact, RED-first) | [mtp-spec-decode.md §3,§5](specs/mtp-spec-decode.md) | `ACTIVE` | `CLAIM-SPEC-GDN-I4`, `CLAIM-SPEC-MTP-I5A` |
-| `SPEC-DFLASH` | Block-diffusion drafter. **READINESS RE-ASSESSED 2026-07-25 (`CLAIM-SPEC-DFLASH-READINESS`, design-only, DONE) against the LANDED MTP machinery (`SPEC-MTP` I1..I7).** Verdict **GREEN, dispatch-ready, NO hardware/oracle/download blocker** (spec [§0](specs/dflash-spec-decode.md)). Refreshed reuse-vs-new map: DFlash gets FREE from landed MTP — the frozen spec-metadata ABI, the greedy rejection sampler (k-general, I3 tested k∈{1,3}), the GDN spec slot path + rollback + mixed spec/non-spec batch (`GdnBlockPagedMixedSpec`/`IndexSelect`/`IndexCopy`, general `num_spec`), the widened-cache-aware conv ops (I5e), the draft-KV layer pattern (`fa_draft`), the I5d/I7 runner verify/propose loop, and **`num_lookahead_tokens=k+1` ALREADY coded** (`speculative.h:91-108` `use_dflash()`); EXTENDS the single I5d-pre `hidden_tap` seam to multi-tap `[T,H×taps]`; builds NEW the `qwen3_dflash` drafter, the project's FIRST non-causal in-block attention primitive, context-KV precompute, `prepare_dflash_inputs`, and the uniform-1+k FULL CG. **k>1 verdict:** the landed rejection + GDN machinery is MECHANICALLY k-general (no `k==1` hardwiring) — DFlash's k=15 blocks need NO mechanism extension, only exercise/validation at scale (D4) + the k+1-slot memory measurement (~2.3 GiB/req 27B GDN state at block-16, the #1 risk, §5). **Checkpoint-fit:** both z-lab drafts EXIST on HF (27B 1.73 GB / 35B 368 MB bf16, DFlashDraftModel) and FIT the 119 GiB pool trivially (drafts NOT yet on dgx — D0 downloads ≤1.73 GB); the active dgx oracle `vllm-oracle-v0.25.0-stage` CONSTRUCTS DFlash (registry `DFlashDraftModel→qwen3_dflash`, speculator dir present) — soft D0 risk = confirm it SERVES DFlash+NVFP4 on sm_121 (non-causal backend; community `AEON-7/vllm-dflash` container proves the combination runs on GB10). W-plan D0-D6 in the spec. **D0+D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`) on the ADVANCED pin `555967922`/vLLM 0.26.0.dev0 — `SPEC-DFLASH` → `ACTIVE`.** D0 UNBLOCKED (vllm#40898 resolved under `VLLM_USE_V2_MODEL_RUNNER=1`): the mixed-attn z-lab 27B draft CONSTRUCTS + the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV, goldens committed); gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF — the k=16 block verify diverges at bf16 near-ties, so NOT the MTP three-way identity). D1 `DF-AUX-TAPS` DONE: `Qwen3_5AuxTaps` + `ModelForwardInput::aux_tap` route to `Qwen3_5{,Dense}Model::ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids` into `[T,H×taps]` (eagle3 `_maybe_add_hidden_state`, aux key L+1); config-gated byte-identical off. Unit gate 598 assertions (independent truncated-model reference, RED-first reversed-concat 384 fail); CUDA 697/697 + compute-sanitizer 0; INERTNESS PROVEN — 27B MTP e2e 9/9 + 27B text SACRED 235/235 byte-identical on the new oracle. **D2 `DF-DRAFT-MODEL` CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D2`, kernel row `KERNEL-ATTN-DFLASH-BLOCK`):** the `qwen3_dflash` draft model (plain 5-layer Qwen3-dense reusing `dense_attn_block.h` ops), the project's FIRST non-causal / bidirectional attention primitive `vt::DFlashBlockAttention` (a SEPARATE op — causal `kAttention`/`kPagedAttention` byte-identical), the fc aux-combine, mask-embed, per-layer SWA/full resolution, and the z-lab loader. CPU gate GREEN (op 12/12 incl. RED non-causal; model forward 95/95 incl. RED full-layer-causal-flip + block isolation + fc RED); existing causal `test_ops_attention` 9/9 + `test_qwen3_forward` 1028 UNCHANGED. **D2 GPU PROMOTION GREEN on dgx (`CLAIM-DFLASH-D2`):** CUDA `-Werror` clean, CUDA==CPU 198412/198412 + compute-sanitizer 0, draft-forward parity vs the REAL vLLM draft (fc rel-L2 0.46%, hidden ≤1.3%, 11 STRICT + 5 near-tie ids), 27B SACRED 235/235 + MTP 9/9 byte-identical — **D2 DONE.** **D3 `DF-DRAFT-KV-PREP` DONE 2026-07-26 (`CLAIM-DFLASH-D3`):** `PrecomputeContextKV` + `PrepareDflashInputs` + `ForwardBlockLogitsWithContext` (reuse the UNCHANGED D2 kernel via [context;block]); GPU numeric-parity `test_qwen3_dflash_kvprep_parity` 61/61 (prepare INTEGER bit-exact vs vLLM's Triton kernel, context-KV K/V rel-L2 0.31%/0.26%, 13 STRICT + 3 near-tie = 16/16), CPU 114/114 RED-proven, inertness 235/235 + 9/9 + D2 37/37 byte-identical. **D4 `DF-ENGINE-INTEGRATION` propose brick + `dflash` config-select CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D4D5`):** `DflashProposeBlock`/`SampleDflashBlockDrafts` (the non-autoregressive whole-block propose composing D3 `ForwardBlockLogitsWithContext` + greedy per-mask argmax, anchor not sampled, `dflash/speculator.py:300-413`) + `ParseSpeculativeConfigJson`/`ResolveDflash` accept `method:"dflash"`. CPU gate `test_dflash_propose` 5/19 GREEN (RED-first anchor-read fails 4/5; brick composes forward+sampler; empty-ctx degenerates to D2; config lookahead k+1). Additive + config-gated ⇒ MTP + non-spec byte-identical BY CONSTRUCTION (`git diff --stat` = new speculator TU + config accept-list + CMake + test, NO runner/model/loader/scheduler edit). **D5 `DF-ENGINE-INTEGRATION` runner-loop LANDED + e2e RUNS on dgx 2026-07-26 (`CLAIM-DFLASH-D5`):** full verify/propose loop wired — loader loads the SEPARATE z-lab draft (`LoadDflashDraft`, host bf16 + target-SHARED bf16 embed/lm_head) via a `--speculative-config` `model` key + `ResolveSpecConfig` dflash branch + `runner.set_dflash_draft`; the verify forward captures the D1 multi-tap (`aux_tap`→`ForwardDeviceMultiTap`) instead of the MTP single tap; `propose_drafts_dflash` ACCUMULATES the per-request combined-feature context (`CombineAuxFeatures(aux_tap)`) across steps and honors the `num_rejected` rollback by appending only the `(T_req−num_rejected)` accepted-prefix features, then runs `DflashProposeBlock` (k=16 GDN-spec exercised first time). **e2e (`test_qwen27_dflash_spec_decode`, 4 prompts×32 tok, our-DFlash-ON vs the committed vLLM-DFlash-ON golden): 2/4 STRICT token-exact (fibonacci, three-laws) + acceptance ~ vLLM on ALL 4 (accepted 19/39/29/25 vs golden 17/39/30/25, deltas +2/0/−1/0 — the MANDATORY dead-drafter-trap condition MET).** The 2 divergences (France tok11 `2972`↔`11751`, 17*23 tok12 `567`↔`488`) are SINGLE bf16 near-tie flips (17*23 RE-CONVERGES after one token = proven near-tie; France cascades from one flip) — the ratified near-tie ROOT the D0 gate-form anticipated, rooted in the D3-documented inline bf16 context-KV recompute envelope (~0.3-1.3% rel-L2), NOT a wiring bug (proven by the 2 exact prompts + near-exact acceptance + a non-trivial shared prefix). Inertness GREEN on this build: SACRED `test_qwen27_paged_engine` 235/235 + MTP `test_qwen27_spec_decode` 9/9 byte-identical; CUDA `-Werror` clean; NO new CUDA kernel (host orchestration reusing D1/D2/D3-sanitized ops). **NOT a clean strict-4/4 pass; STRICT 4/4 token-identity + the speed A/B = D6 (the persistent paged draft-KV bit-matching vLLM's fused context-KV projections + the uniform-1+k FULL CG).** Row STAYS `ACTIVE` (correctness at the ratified near-tie envelope; D6 remains) **D6 2026-07-27 (`CLAIM-DFLASH-D6`) — c1 SPEED A/B DONE + STRICT-irreducibility RCA + CG feasibility (records-only, NO source code):** (1) **c1 speed A/B** (`examples/vllm-bench` at `361189a7`, 8 prose+code prompts×256 tok greedy c1, 2 reps): our DFlash-ON = **2.50x TPOT (40.4 vs 101.2 ms) / 2.48x output-tput (24.4 vs 9.86 tok/s)** over our OFF, acceptance 0.22 (3.56/16), rep-stable <1.5%; `benchmark_binding=true`. vs vLLM-DFlash-ON graphed (same workload): vLLM-DFlash-ON graphed = 28.5 tok/s / 35.1 ms TPOT / acceptance_len 4.30 (same 8 prompts, `VLLM_USE_V2_MODEL_RUNNER=1`, mm-off, gpu_util 0.30), so OURS IS ~14% BELOW vLLM-DFlash-ON on output throughput (24.4 vs 28.5 tok/s) - both ~on-par at spec-OFF (9.86 vs 9.83 tok/s), but vLLM extracts a larger DFlash speedup (2.90x vs our 2.47x) because its draft step is fully device-resident + CUDA-graphed (ours host-orchestrates 13 downloads/step) + slightly higher acceptance (~4.3 vs ~3.6 draft tokens/step). The DONE speed bar (ours >= vLLM) is NOT met; closing it = the device-resident draft rewrite + FULL CG (D6 part 2). (2) **STRICT-4/4 proven bf16-IRREDUCIBLE** — the draft KV cache is bf16 not fp8 (`torch_utils.py:398` `auto`→model dtype; the D0 "fp8-KV" was the backend name, not the KV storage dtype), the D3 golden already compares pre-storage bf16 (residual K 0.31%/V 0.26% = sub-ULP kernel noise), and a fused multi-layer KV GEMM is per-element invariant to our per-layer GEMMs ⇒ bit-exact needs vLLM's exact kernels ⇒ the ratified near-tie gate is the FINAL correctness form (no fused-KV code landed). (3) **FULL CG BLOCKED** on a device-resident draft-path rewrite (the D5 path does 13 device→host downloads/step + host `[context;block]` interleaving) — the remaining throughput-parity increment (the perf form of persistent-paged-KV + the graph). Inertness by construction (the gated binary is the D5 binary; SACRED 235/235 + MTP 9/9 stand). Evidence tool `scripts/spec/vllm_dflash_timing.py`. **D7 2026-07-27 (`CLAIM-DFLASH-D7`) — within-step draft forward made DEVICE-RESIDENT (source-owning): `PrecomputeContextKVDevice` keeps per-layer K/V on device; `ForwardBlockLogitsWithContext` builds [context;block] with `vt::IndexCopy`/`IndexSelect` (removes ~30 D→H `Download`s/step). BIT-IDENTICAL (identity bf16↔f32 round-trips replaced) — e2e `test_qwen27_dflash_spec_decode` 27/27 SAME tokens (2/4 STRICT + 2/4 near-tie, acceptance 19/39/29/25), SACRED 235/235 + MTP 9/9, CUDA `-Werror` clean, compute-sanitizer 0 (198412). But the direct old-vs-new A/B = +2.0% output-tput (IN-NOISE) ⇒ D6's "downloads = the ~14% gap" REFUTED by measurement; ours 19.68 tok/s STILL ~33% BELOW vLLM-DFlash-ON 29.2 tok/s (reconstructed 8-prompt set, more prose-heavy); OFF parity our 9.97 ≥ vLLM 9.66. Residual re-attributed: acceptance (ours 2.49 vs vLLM ~3.13 accepted draft-tok/step, bf16-irreducible) + per-step context-KV RECOMPUTE (O(context²), needs the cross-step persistent paged draft-KV store) + eager-vs-graphed. SPEED BAR NOT met; SPEC-DFLASH stays `ACTIVE`; next = persistent paged draft-KV store → then FULL CG. **D9 2026-07-27 (`CLAIM-DFLASH-D9`) — PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7% throughput, 0.69×→0.917×); D8 acceptance-ceiling REFUTED; residual = FULL CG ONLY:** `qwen3_dflash.cpp` `AppendContextKVHost` (project ONLY newly-accepted rows → per-layer bf16 K/V, append to `PrecomputedContextKV`) + `ForwardBlockLogitsWithPrecomputedKV` (upload the persistent store, NO re-projection) share the core `ForwardWithCtxKVDev` with the old recompute; `runner.cpp::propose_drafts_dflash` swaps the O(context²) per-step recompute (`dflash_ctx_feats_`) for an append-only per-request `dflash_kv_store_` (rollback=don't-append). NO new CUDA kernel; config-gated. BIT-IDENTICAL: CPU `test_dflash_propose` two new D9 cases = exact float equality vs full recompute; GPU e2e `test_qwen27_dflash_spec_decode` **27/27 SAME tokens** (acceptance 19/39/29/25, same divergences France@11/17×23@12); SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean. **A/B (c1, 8 prose+code×256 tok input-len 512, 2 reps <0.1%, `benchmark_binding=true`):** ours-ON **25.75 tok/s** (was D8 20.99, +22.7%) / 38.40 ms TPOT / acc **3.68/step** vs vLLM-ON graphed **28.09** / 35.60 / acc 3.31 = **0.917×** (~8% below, was 0.69×). **Part 1 same-trajectory:** on the 2 token-identical-trajectory prompts ours per-step acceptance == vLLM's EXACTLY (fibonacci 7.80/7.80, three-laws 3.571/3.571, ratio 1.00) AND on the A/B ours acceptance (3.68) is HIGHER than vLLM's (3.31) ⇒ D8's 0.80–0.85× "bf16 acceptance ceiling" is a trajectory-divergence CONFOUND, REFUTED. Residual (~8%) = eager-vs-graphed ONLY (ours ON/OFF 2.60× vs vLLM 2.91×, OFF at parity, recompute eliminated, acceptance higher) — NOT an irreducible ceiling; the FULL uniform-(1+k) CG (device paged-KV store + paged attn, new-CUDA multi-file) is the SOLE un-landed increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM; residual isolated to FULL CG). **D12 2026-07-27 (`CLAIM-DFLASH-D12`) — A-wire + Part B LANDED + GPU-gated; Part C (capture) remaining; 0.917×:** A-wire makes the D11 Part-A device store the PRODUCTION path (`runner.{h,cpp}` `dflash_kv_store_`→`shared_ptr<DflashDeviceKVStore>`, `MakeDeviceKVStore`/`AppendContextKVDevice`/`ForwardBlockLogitsWithDeviceKV`; GPU-gated e2e `test_qwen27_dflash_spec_decode` 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical, `-Werror` clean). Part B adds `vt::DFlashPagedBlockAttention` (`OpId::kDFlashPagedBlockAttention`), the capture-safe paged kernel with EVERY metadata input a persistent DEVICE tensor and NO function-local host `cu_seqlens` upload (fixes the `cuda_ops.cu:1277-1280` capture-UAF class), gated CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` `test_ops_dflash_paged_block_attn` 795648/795648 + compute-sanitizer 0. Speed 0.917× (A-wire eager + Part B not yet wired into the forward); `benchmark_binding=false`. Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B) is the SOLE remaining piece; if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE. Stays `ACTIVE`. **D13 2026-07-27 (`CLAIM-DFLASH-D13`) — Part C LANDED + GPU-GATED; capture-correctness PROVEN; c1 throughput NEAR-PARITY (ours 0.978x, ~2% below vLLM); gap CLOSED 0.917x→0.978x; STAYS `ACTIVE` (≥vLLM bar not yet met):** single-file additive change (`qwen3_dflash.cpp` +368/-58). (C.1) `DflashDeviceKVStore` → fixed-capacity PAGED cache (per-layer pool `[max_pages,16,Hkv,Dh]` + identity `block_table` + `seq_lens`; append = `vt::IndexCopy` scatter at slot==abs-pos, bit-identical to the D9/D11 store). (C.2) `ForwardPagedBody` runs the (1+k) block through the D12 `vt::DFlashPagedBlockAttention` reading the paged store (no `[context;block]` materialization, no function-local host uploads); runner P==1 propose routes through it, P>1 bit-identical materialized fallback. (C.3) per-request CUDA GRAPH over the paged draft step (warm-in-step repopulates the shared pool free-list right before `BeginCapture` — the fix for a `cudaMalloc`-in-capture `Get` miss from the intervening 27B target forward — then `BeginCapture → ForwardPagedBody → EndCaptureGraph`, replay with growing context entering only via in-place `seq_lens`). **Capture-correctness (MANDATORY): `test_qwen27_dflash_spec_decode` 27/27 with the graph (VT_DFLASH_GRAPH=1) BIT-IDENTICAL to eager (=0)** — same divergence tokens (France@11 got[…2972…], 17×23@12 got[…567…]), same acceptance 19/39/29/25 as D5/D7/D9/D12; graph ENGAGED (5 captures C=2048/5/4/15/6, 32+ replays); the token-diff is the capture-safety proof ([[cudagraph-capture-bakes-stack-addresses]]). **c1 A/B (one flock series, cold rep discarded, 8 prompts×256 tok):** our OFF 10.24 / our ON eager-paged 28.65 (28.69,28.61) / **our ON GRAPHED 28.70 (28.70,28.70), TPOT 34.40** / vLLM-ON graphed steady-state 29.35 (tight 3-rep 29.33/29.37/29.33, TPOT 34.07, acc_len 4.44); D9's 28.09 was a colder cross-session outlier — **NEAR-PARITY: ours 0.978× (~2% below) on the rigorous same-session band** (across sessions ours 28.70 falls inside vLLM's observed 28.09–29.37 range). ON/OFF 2.80× (vLLM ~2.98×), our OFF ≥ vLLM OFF. Per the acceptance rule ("below on any axis = an open gap; near-parity is NOT met"), the ≥vLLM bar is NOT met; STAYS `ACTIVE`. Residual (data-grounded): NOT acceptance (ours realized ~3.68 accepted draft-tok/step > vLLM's 3.44) and NOT launch/graph (both graphed, CG neutral) — per-step COMPUTE (~2% slower target-step); next lever = nsys both draft steps (`--cuda-graph-trace=node`), no premature ceiling. **ATTRIBUTION (supersedes D9):** the CUDA graph is perf-NEUTRAL (+0.3%); the ACTUAL lever was the paged context read (C.1/C.2) removing the D9/D12 per-layer `[context;block]` `IndexCopy` materialization of the whole growing context (25.75 D9 → 28.65 eager-paged, +11%) — the roadmap's "the full CG closes the gap" premise is corrected by measurement. Inertness VERIFIED on the capture binary: SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel (D12 paged kernel already memcheck-0 795648), `check-device-leakage` not increased (paged path REMOVES the materialized-buffer allocs + host uploads). `benchmark_binding=true`. Correctness-complete (ratified near-tie); throughput NEAR-PARITY (0.978×, ~2% residual) ⇒ STAYS `ACTIVE` (the capture-correctness gate is MET; the ≥vLLM speed bar is the sole remaining item, a ~2% per-step-compute residual for an nsys). Anchors: `src/vllm/model_executor/models/qwen3_dflash.cpp` (`DflashDeviceKVStore` paged store, `ForwardPagedBody`, the per-request graph in `ForwardBlockLogitsWithDeviceKV`). **D14 2026-07-27 (`CLAIM-DFLASH-D14`) — SPEED GATE MET → SPEC-DFLASH `DONE`:** an nsys (`--cuda-graph-trace=node`) of the graphed spec-on step attributed the D13 ~2% residual to the from-scratch `DFlashPagedBlockAttentionKernel` draft attention (242.9 ms = 1.8% of GPU time, median ~460 us/call over context C~500-640, vs vLLM's fused flash draft-attn ~0.15%; BOTH engines run identical `cutlass_80_wmma` for the draft bf16 GEMMs, so the GEMMs were NOT the gap). Ported it to a WARP-scoped online-softmax variant `DFlashPagedBlockAttentionWarpKernel` (mirrors the shipped `AttentionWarpKernel`: one warp per (block-query,head), `__shfl_xor` butterfly reduction, register accumulator, NO `__syncthreads` storm; SAME paged/block combined-index read + causal/SWA mask + GQA; default ON, `VT_DFLASH_ATTN_BLOCK=1` keeps the bit-identical D12/D13 block kernel for A/B). Draft attn 242.9 → 77.9 ms (3.1x); our-ON c1 28.60 → 29.32 tok/s (+2.5%). **FINAL same-session 3-rep A/B (8 prompts×256 tok, cold leg discarded): our-ON graphed 29.42/29.27/29.32 (med 29.32) vs vLLM-ON graphed 29.240/29.247/29.233 (med 29.240) — our WORST rep (29.27) > vLLM's BEST (29.247), NON-OVERLAPPING bands, 1.003× ⇒ the ≥vLLM speed gate is MET.** Correctness UNCHANGED (output is exact by spec-decode construction — the target verify is untouched, only which draft proposals are accepted can shift): e2e `test_qwen27_dflash_spec_decode` 27/27 with graph==eager BIT-IDENTICAL, acceptance 19/39/29/25 unchanged (draft accepted 1629 identical warp-vs-block across the whole A/B set), 2/4 STRICT (France@11, 17×23@12 unchanged); CUDA==CPU `test_ops_dflash_paged_block_attn` 795648/795648 (warp within the f32 1e-4 / bf16 3e-2 envelope) + compute-sanitizer 0. Inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean; `check-device-leakage` not increased. `benchmark_binding=true`. Block-diffusion drafting is now correctness-complete (ratified near-tie) AND at/above vLLM throughput — this was the roadmap's FINAL open speed item. Anchors: `src/vt/cuda/cuda_ops.cu` (`DFlashPagedBlockAttentionWarpKernel` + `UseDflashAttnBlockKernel`; the D12 block kernel retained as the `VT_DFLASH_ATTN_BLOCK=1` reference). | T1 | `vllm/v1/worker/gpu/spec_decode/dflash/speculator.py`; `vllm/model_executor/models/qwen3_dflash.py`; `vllm/model_executor/models/interfaces.py:1382` (aux value); `eagle3_utils.py:41-56` (+1 shift) | `include/vllm/model_executor/models/qwen3_5.h` (`Qwen3_5AuxTaps`, `ForwardDeviceMultiTap`); `qwen3_5_dense.h`; `model_registry.h` (`aux_tap`); `src/vllm/model_executor/models/qwen3_5.cpp` (`MaybeCaptureAuxTap`/`ValidateAuxTapLayerIds`/`ForwardDeviceMultiTap`); `qwen3_5_moe.cpp`+`qwen3_5_dense.cpp` (routing); D2/D3 `include/vllm/model_executor/models/qwen3_dflash.h` + `src/vllm/model_executor/models/qwen3_dflash{,_weights}.cpp`; D4 `include/vllm/v1/worker/gpu/spec_decode/dflash/speculator.h` + `src/vllm/v1/worker/gpu/spec_decode/dflash/speculator.cpp` (`DflashProposeBlock`/`SampleDflashBlockDrafts`); D5 `src/vllm/entrypoints/model_loader.cpp` (`LoadDflashDraft`/`DflashDraft`) + `include/vllm/entrypoints/model_loader.h`; D5 `src/vllm/v1/worker/gpu/runner.cpp` (`set_dflash_draft`/`propose_drafts_dflash`/aux-tap capture) + `include/vllm/v1/worker/gpu/runner.h`; `src/vllm/config/speculative.cpp` + `include/vllm/config/speculative.h` (`ResolveDflash` + `dflash`/`model` parse); D14 warp kernel [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu#L1433) | `tests/vllm/models/test_qwen27_paged_forward.cpp` (multi-tap 598); `tests/vt/test_ops_dflash_block_attn.cpp`; `tests/vllm/models/test_qwen3_dflash_forward.cpp`; `tests/vllm/v1/spec_decode/test_dflash_kvprep.cpp`; `tests/parity/test_qwen3_dflash_{draft,kvprep}_parity.cpp`; D4 `tests/vllm/v1/spec_decode/test_dflash_propose.cpp` (5/19, RED-first); D5 `tests/parity/test_qwen27_dflash_spec_decode.cpp` (e2e 27/27, 2/4 strict + acceptance~vLLM); `scripts/spec/d{0,2,3}_dflash_*.py`; `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`; D6 `scripts/spec/vllm_dflash_timing.py` (vLLM-DFlash c1 timing); D7 device-resident `src/vllm/model_executor/models/qwen3_dflash.cpp` (`PrecomputeContextKVDevice` + `ForwardBlockLogitsWithContext` via `vt::IndexCopy`/`IndexSelect`); D9 persistent paged draft-KV `qwen3_dflash.{h,cpp}` (`AppendContextKVHost`/`ForwardBlockLogitsWithPrecomputedKV`/`ForwardWithCtxKVDev`/`PrecomputedContextKV`) + `runner.{h,cpp}` (`dflash_kv_store_`/`propose_drafts_dflash`) + `tests/vllm/v1/spec_decode/test_dflash_propose.cpp` (2 D9 bit-identity cases); D12 A-wire `runner.{h,cpp}` (device store as production path) + D12 Part B `include/vt/ops.h`/`src/vt/ops.cpp`/`src/vt/cpu/cpu_ops.cpp`/`src/vt/cuda/cuda_ops.cu` (`kDFlashPagedBlockAttention`) + `tests/vt/test_ops_dflash_paged_block_attn.cpp` (CPU==CUDA + cross-check, 795648/795648 + sanitizer-0); D13 `src/vllm/model_executor/models/qwen3_dflash.cpp` (fixed-capacity paged `DflashDeviceKVStore` + `ForwardPagedBody` + the per-request draft-step CUDA graph in `ForwardBlockLogitsWithDeviceKV`); D14 [test_ops_dflash_paged_block_attn](../tests/vt/test_ops_dflash_paged_block_attn.cpp#L79) + [ledger](parity-ledger.md#L738) | [dflash-spec-decode.md](specs/dflash-spec-decode.md) | `DONE` | `164453a2` |
+| `SPEC-DFLASH` | Block-diffusion drafter. **READINESS RE-ASSESSED 2026-07-25 (`CLAIM-SPEC-DFLASH-READINESS`, design-only, DONE) against the LANDED MTP machinery (`SPEC-MTP` I1..I7).** Verdict **GREEN, dispatch-ready, NO hardware/oracle/download blocker** (spec [§0](specs/dflash-spec-decode.md)). Refreshed reuse-vs-new map: DFlash gets FREE from landed MTP — the frozen spec-metadata ABI, the greedy rejection sampler (k-general, I3 tested k∈{1,3}), the GDN spec slot path + rollback + mixed spec/non-spec batch (`GdnBlockPagedMixedSpec`/`IndexSelect`/`IndexCopy`, general `num_spec`), the widened-cache-aware conv ops (I5e), the draft-KV layer pattern (`fa_draft`), the I5d/I7 runner verify/propose loop, and **`num_lookahead_tokens=k+1` ALREADY coded** (`speculative.h:91-108` `use_dflash()`); EXTENDS the single I5d-pre `hidden_tap` seam to multi-tap `[T,H×taps]`; builds NEW the `qwen3_dflash` drafter, the project's FIRST non-causal in-block attention primitive, context-KV precompute, `prepare_dflash_inputs`, and the uniform-1+k FULL CG. **k>1 verdict:** the landed rejection + GDN machinery is MECHANICALLY k-general (no `k==1` hardwiring) — DFlash's k=15 blocks need NO mechanism extension, only exercise/validation at scale (D4) + the k+1-slot memory measurement (~2.3 GiB/req 27B GDN state at block-16, the #1 risk, §5). **Checkpoint-fit:** both z-lab drafts EXIST on HF (27B 1.73 GB / 35B 368 MB bf16, DFlashDraftModel) and FIT the 119 GiB pool trivially (drafts NOT yet on dgx — D0 downloads ≤1.73 GB); the active dgx oracle `vllm-oracle-v0.25.0-stage` CONSTRUCTS DFlash (registry `DFlashDraftModel→qwen3_dflash`, speculator dir present) — soft D0 risk = confirm it SERVES DFlash+NVFP4 on sm_121 (non-causal backend; community `AEON-7/vllm-dflash` container proves the combination runs on GB10). W-plan D0-D6 in the spec. **D0+D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`) on the ADVANCED pin `555967922`/vLLM 0.26.0.dev0 — `SPEC-DFLASH` → `ACTIVE`.** D0 UNBLOCKED (vllm#40898 resolved under `VLLM_USE_V2_MODEL_RUNNER=1`): the mixed-attn z-lab 27B draft CONSTRUCTS + the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV, goldens committed); gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF — the k=16 block verify diverges at bf16 near-ties, so NOT the MTP three-way identity). D1 `DF-AUX-TAPS` DONE: `Qwen3_5AuxTaps` + `ModelForwardInput::aux_tap` route to `Qwen3_5{,Dense}Model::ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids` into `[T,H×taps]` (eagle3 `_maybe_add_hidden_state`, aux key L+1); config-gated byte-identical off. Unit gate 598 assertions (independent truncated-model reference, RED-first reversed-concat 384 fail); CUDA 697/697 + compute-sanitizer 0; INERTNESS PROVEN — 27B MTP e2e 9/9 + 27B text SACRED 235/235 byte-identical on the new oracle. **D2 `DF-DRAFT-MODEL` CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D2`, kernel row `KERNEL-ATTN-DFLASH-BLOCK`):** the `qwen3_dflash` draft model (plain 5-layer Qwen3-dense reusing `dense_attn_block.h` ops), the project's FIRST non-causal / bidirectional attention primitive `vt::DFlashBlockAttention` (a SEPARATE op — causal `kAttention`/`kPagedAttention` byte-identical), the fc aux-combine, mask-embed, per-layer SWA/full resolution, and the z-lab loader. CPU gate GREEN (op 12/12 incl. RED non-causal; model forward 95/95 incl. RED full-layer-causal-flip + block isolation + fc RED); existing causal `test_ops_attention` 9/9 + `test_qwen3_forward` 1028 UNCHANGED. **D2 GPU PROMOTION GREEN on dgx (`CLAIM-DFLASH-D2`):** CUDA `-Werror` clean, CUDA==CPU 198412/198412 + compute-sanitizer 0, draft-forward parity vs the REAL vLLM draft (fc rel-L2 0.46%, hidden ≤1.3%, 11 STRICT + 5 near-tie ids), 27B SACRED 235/235 + MTP 9/9 byte-identical — **D2 DONE.** **D3 `DF-DRAFT-KV-PREP` DONE 2026-07-26 (`CLAIM-DFLASH-D3`):** `PrecomputeContextKV` + `PrepareDflashInputs` + `ForwardBlockLogitsWithContext` (reuse the UNCHANGED D2 kernel via [context;block]); GPU numeric-parity `test_qwen3_dflash_kvprep_parity` 61/61 (prepare INTEGER bit-exact vs vLLM's Triton kernel, context-KV K/V rel-L2 0.31%/0.26%, 13 STRICT + 3 near-tie = 16/16), CPU 114/114 RED-proven, inertness 235/235 + 9/9 + D2 37/37 byte-identical. **D4 `DF-ENGINE-INTEGRATION` propose brick + `dflash` config-select CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D4D5`):** `DflashProposeBlock`/`SampleDflashBlockDrafts` (the non-autoregressive whole-block propose composing D3 `ForwardBlockLogitsWithContext` + greedy per-mask argmax, anchor not sampled, `dflash/speculator.py:300-413`) + `ParseSpeculativeConfigJson`/`ResolveDflash` accept `method:"dflash"`. CPU gate `test_dflash_propose` 5/19 GREEN (RED-first anchor-read fails 4/5; brick composes forward+sampler; empty-ctx degenerates to D2; config lookahead k+1). Additive + config-gated ⇒ MTP + non-spec byte-identical BY CONSTRUCTION (`git diff --stat` = new speculator TU + config accept-list + CMake + test, NO runner/model/loader/scheduler edit). **D5 `DF-ENGINE-INTEGRATION` runner-loop LANDED + e2e RUNS on dgx 2026-07-26 (`CLAIM-DFLASH-D5`):** full verify/propose loop wired — loader loads the SEPARATE z-lab draft (`LoadDflashDraft`, host bf16 + target-SHARED bf16 embed/lm_head) via a `--speculative-config` `model` key + `ResolveSpecConfig` dflash branch + `runner.set_dflash_draft`; the verify forward captures the D1 multi-tap (`aux_tap`→`ForwardDeviceMultiTap`) instead of the MTP single tap; `propose_drafts_dflash` ACCUMULATES the per-request combined-feature context (`CombineAuxFeatures(aux_tap)`) across steps and honors the `num_rejected` rollback by appending only the `(T_req−num_rejected)` accepted-prefix features, then runs `DflashProposeBlock` (k=16 GDN-spec exercised first time). **e2e (`test_qwen27_dflash_spec_decode`, 4 prompts×32 tok, our-DFlash-ON vs the committed vLLM-DFlash-ON golden): 2/4 STRICT token-exact (fibonacci, three-laws) + acceptance ~ vLLM on ALL 4 (accepted 19/39/29/25 vs golden 17/39/30/25, deltas +2/0/−1/0 — the MANDATORY dead-drafter-trap condition MET).** The 2 divergences (France tok11 `2972`↔`11751`, 17*23 tok12 `567`↔`488`) are SINGLE bf16 near-tie flips (17*23 RE-CONVERGES after one token = proven near-tie; France cascades from one flip) — the ratified near-tie ROOT the D0 gate-form anticipated, rooted in the D3-documented inline bf16 context-KV recompute envelope (~0.3-1.3% rel-L2), NOT a wiring bug (proven by the 2 exact prompts + near-exact acceptance + a non-trivial shared prefix). Inertness GREEN on this build: SACRED `test_qwen27_paged_engine` 235/235 + MTP `test_qwen27_spec_decode` 9/9 byte-identical; CUDA `-Werror` clean; NO new CUDA kernel (host orchestration reusing D1/D2/D3-sanitized ops). **NOT a clean strict-4/4 pass; STRICT 4/4 token-identity + the speed A/B = D6 (the persistent paged draft-KV bit-matching vLLM's fused context-KV projections + the uniform-1+k FULL CG).** Row STAYS `ACTIVE` (correctness at the ratified near-tie envelope; D6 remains) **D6 2026-07-27 (`CLAIM-DFLASH-D6`) — c1 SPEED A/B DONE + STRICT-irreducibility RCA + CG feasibility (records-only, NO source code):** (1) **c1 speed A/B** (`examples/vllm-bench` at `361189a7`, 8 prose+code prompts×256 tok greedy c1, 2 reps): our DFlash-ON = **2.50x TPOT (40.4 vs 101.2 ms) / 2.48x output-tput (24.4 vs 9.86 tok/s)** over our OFF, acceptance 0.22 (3.56/16), rep-stable <1.5%; `benchmark_binding=true`. vs vLLM-DFlash-ON graphed (same workload): vLLM-DFlash-ON graphed = 28.5 tok/s / 35.1 ms TPOT / acceptance_len 4.30 (same 8 prompts, `VLLM_USE_V2_MODEL_RUNNER=1`, mm-off, gpu_util 0.30), so OURS IS ~14% BELOW vLLM-DFlash-ON on output throughput (24.4 vs 28.5 tok/s) - both ~on-par at spec-OFF (9.86 vs 9.83 tok/s), but vLLM extracts a larger DFlash speedup (2.90x vs our 2.47x) because its draft step is fully device-resident + CUDA-graphed (ours host-orchestrates 13 downloads/step) + slightly higher acceptance (~4.3 vs ~3.6 draft tokens/step). The DONE speed bar (ours >= vLLM) is NOT met; closing it = the device-resident draft rewrite + FULL CG (D6 part 2). (2) **STRICT-4/4 proven bf16-IRREDUCIBLE** — the draft KV cache is bf16 not fp8 (`torch_utils.py:398` `auto`→model dtype; the D0 "fp8-KV" was the backend name, not the KV storage dtype), the D3 golden already compares pre-storage bf16 (residual K 0.31%/V 0.26% = sub-ULP kernel noise), and a fused multi-layer KV GEMM is per-element invariant to our per-layer GEMMs ⇒ bit-exact needs vLLM's exact kernels ⇒ the ratified near-tie gate is the FINAL correctness form (no fused-KV code landed). (3) **FULL CG BLOCKED** on a device-resident draft-path rewrite (the D5 path does 13 device→host downloads/step + host `[context;block]` interleaving) — the remaining throughput-parity increment (the perf form of persistent-paged-KV + the graph). Inertness by construction (the gated binary is the D5 binary; SACRED 235/235 + MTP 9/9 stand). Evidence tool `scripts/spec/vllm_dflash_timing.py`. **D7 2026-07-27 (`CLAIM-DFLASH-D7`) — within-step draft forward made DEVICE-RESIDENT (source-owning): `PrecomputeContextKVDevice` keeps per-layer K/V on device; `ForwardBlockLogitsWithContext` builds [context;block] with `vt::IndexCopy`/`IndexSelect` (removes ~30 D→H `Download`s/step). BIT-IDENTICAL (identity bf16↔f32 round-trips replaced) — e2e `test_qwen27_dflash_spec_decode` 27/27 SAME tokens (2/4 STRICT + 2/4 near-tie, acceptance 19/39/29/25), SACRED 235/235 + MTP 9/9, CUDA `-Werror` clean, compute-sanitizer 0 (198412). But the direct old-vs-new A/B = +2.0% output-tput (IN-NOISE) ⇒ D6's "downloads = the ~14% gap" REFUTED by measurement; ours 19.68 tok/s STILL ~33% BELOW vLLM-DFlash-ON 29.2 tok/s (reconstructed 8-prompt set, more prose-heavy); OFF parity our 9.97 ≥ vLLM 9.66. Residual re-attributed: acceptance (ours 2.49 vs vLLM ~3.13 accepted draft-tok/step, bf16-irreducible) + per-step context-KV RECOMPUTE (O(context²), needs the cross-step persistent paged draft-KV store) + eager-vs-graphed. SPEED BAR NOT met; SPEC-DFLASH stays `ACTIVE`; next = persistent paged draft-KV store → then FULL CG. **D9 2026-07-27 (`CLAIM-DFLASH-D9`) — PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7% throughput, 0.69×→0.917×); D8 acceptance-ceiling REFUTED; residual = FULL CG ONLY:** `qwen3_dflash.cpp` `AppendContextKVHost` (project ONLY newly-accepted rows → per-layer bf16 K/V, append to `PrecomputedContextKV`) + `ForwardBlockLogitsWithPrecomputedKV` (upload the persistent store, NO re-projection) share the core `ForwardWithCtxKVDev` with the old recompute; `runner.cpp::propose_drafts_dflash` swaps the O(context²) per-step recompute (`dflash_ctx_feats_`) for an append-only per-request `dflash_kv_store_` (rollback=don't-append). NO new CUDA kernel; config-gated. BIT-IDENTICAL: CPU `test_dflash_propose` two new D9 cases = exact float equality vs full recompute; GPU e2e `test_qwen27_dflash_spec_decode` **27/27 SAME tokens** (acceptance 19/39/29/25, same divergences France@11/17×23@12); SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean. **A/B (c1, 8 prose+code×256 tok input-len 512, 2 reps <0.1%, `benchmark_binding=true`):** ours-ON **25.75 tok/s** (was D8 20.99, +22.7%) / 38.40 ms TPOT / acc **3.68/step** vs vLLM-ON graphed **28.09** / 35.60 / acc 3.31 = **0.917×** (~8% below, was 0.69×). **Part 1 same-trajectory:** on the 2 token-identical-trajectory prompts ours per-step acceptance == vLLM's EXACTLY (fibonacci 7.80/7.80, three-laws 3.571/3.571, ratio 1.00) AND on the A/B ours acceptance (3.68) is HIGHER than vLLM's (3.31) ⇒ D8's 0.80–0.85× "bf16 acceptance ceiling" is a trajectory-divergence CONFOUND, REFUTED. Residual (~8%) = eager-vs-graphed ONLY (ours ON/OFF 2.60× vs vLLM 2.91×, OFF at parity, recompute eliminated, acceptance higher) — NOT an irreducible ceiling; the FULL uniform-(1+k) CG (device paged-KV store + paged attn, new-CUDA multi-file) is the SOLE un-landed increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM; residual isolated to FULL CG). **D12 2026-07-27 (`CLAIM-DFLASH-D12`) — A-wire + Part B LANDED + GPU-gated; Part C (capture) remaining; 0.917×:** A-wire makes the D11 Part-A device store the PRODUCTION path (`runner.{h,cpp}` `dflash_kv_store_`→`shared_ptr<DflashDeviceKVStore>`, `MakeDeviceKVStore`/`AppendContextKVDevice`/`ForwardBlockLogitsWithDeviceKV`; GPU-gated e2e `test_qwen27_dflash_spec_decode` 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical, `-Werror` clean). Part B adds `vt::DFlashPagedBlockAttention` (`OpId::kDFlashPagedBlockAttention`), the capture-safe paged kernel with EVERY metadata input a persistent DEVICE tensor and NO function-local host `cu_seqlens` upload (fixes the `cuda_ops.cu:1277-1280` capture-UAF class), gated CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` `test_ops_dflash_paged_block_attn` 795648/795648 + compute-sanitizer 0. Speed 0.917× (A-wire eager + Part B not yet wired into the forward); `benchmark_binding=false`. Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B) is the SOLE remaining piece; if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE. Stays `ACTIVE`. **D13 2026-07-27 (`CLAIM-DFLASH-D13`) — Part C LANDED + GPU-GATED; capture-correctness PROVEN; c1 throughput NEAR-PARITY (ours 0.978x, ~2% below vLLM); gap CLOSED 0.917x→0.978x; STAYS `ACTIVE` (≥vLLM bar not yet met):** single-file additive change (`qwen3_dflash.cpp` +368/-58). (C.1) `DflashDeviceKVStore` → fixed-capacity PAGED cache (per-layer pool `[max_pages,16,Hkv,Dh]` + identity `block_table` + `seq_lens`; append = `vt::IndexCopy` scatter at slot==abs-pos, bit-identical to the D9/D11 store). (C.2) `ForwardPagedBody` runs the (1+k) block through the D12 `vt::DFlashPagedBlockAttention` reading the paged store (no `[context;block]` materialization, no function-local host uploads); runner P==1 propose routes through it, P>1 bit-identical materialized fallback. (C.3) per-request CUDA GRAPH over the paged draft step (warm-in-step repopulates the shared pool free-list right before `BeginCapture` — the fix for a `cudaMalloc`-in-capture `Get` miss from the intervening 27B target forward — then `BeginCapture → ForwardPagedBody → EndCaptureGraph`, replay with growing context entering only via in-place `seq_lens`). **Capture-correctness (MANDATORY): `test_qwen27_dflash_spec_decode` 27/27 with the graph (VT_DFLASH_GRAPH=1) BIT-IDENTICAL to eager (=0)** — same divergence tokens (France@11 got[…2972…], 17×23@12 got[…567…]), same acceptance 19/39/29/25 as D5/D7/D9/D12; graph ENGAGED (5 captures C=2048/5/4/15/6, 32+ replays); the token-diff is the capture-safety proof ([[cudagraph-capture-bakes-stack-addresses]]). **c1 A/B (one flock series, cold rep discarded, 8 prompts×256 tok):** our OFF 10.24 / our ON eager-paged 28.65 (28.69,28.61) / **our ON GRAPHED 28.70 (28.70,28.70), TPOT 34.40** / vLLM-ON graphed steady-state 29.35 (tight 3-rep 29.33/29.37/29.33, TPOT 34.07, acc_len 4.44); D9's 28.09 was a colder cross-session outlier — **NEAR-PARITY: ours 0.978× (~2% below) on the rigorous same-session band** (across sessions ours 28.70 falls inside vLLM's observed 28.09–29.37 range). ON/OFF 2.80× (vLLM ~2.98×), our OFF ≥ vLLM OFF. Per the acceptance rule ("below on any axis = an open gap; near-parity is NOT met"), the ≥vLLM bar is NOT met; STAYS `ACTIVE`. Residual (data-grounded): NOT acceptance (ours realized ~3.68 accepted draft-tok/step > vLLM's 3.44) and NOT launch/graph (both graphed, CG neutral) — per-step COMPUTE (~2% slower target-step); next lever = nsys both draft steps (`--cuda-graph-trace=node`), no premature ceiling. **ATTRIBUTION (supersedes D9):** the CUDA graph is perf-NEUTRAL (+0.3%); the ACTUAL lever was the paged context read (C.1/C.2) removing the D9/D12 per-layer `[context;block]` `IndexCopy` materialization of the whole growing context (25.75 D9 → 28.65 eager-paged, +11%) — the roadmap's "the full CG closes the gap" premise is corrected by measurement. Inertness VERIFIED on the capture binary: SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel (D12 paged kernel already memcheck-0 795648), `check-device-leakage` not increased (paged path REMOVES the materialized-buffer allocs + host uploads). `benchmark_binding=true`. Correctness-complete (ratified near-tie); throughput NEAR-PARITY (0.978×, ~2% residual) ⇒ STAYS `ACTIVE` (the capture-correctness gate is MET; the ≥vLLM speed bar is the sole remaining item, a ~2% per-step-compute residual for an nsys). Anchors: `src/vllm/model_executor/models/qwen3_dflash.cpp` (`DflashDeviceKVStore` paged store, `ForwardPagedBody`, the per-request graph in `ForwardBlockLogitsWithDeviceKV`). **D14 2026-07-27 (`CLAIM-DFLASH-D14`) — SPEED GATE MET → SPEC-DFLASH `DONE`:** an nsys (`--cuda-graph-trace=node`) of the graphed spec-on step attributed the D13 ~2% residual to the from-scratch `DFlashPagedBlockAttentionKernel` draft attention (242.9 ms = 1.8% of GPU time, median ~460 us/call over context C~500-640, vs vLLM's fused flash draft-attn ~0.15%; BOTH engines run identical `cutlass_80_wmma` for the draft bf16 GEMMs, so the GEMMs were NOT the gap). Ported it to a WARP-scoped online-softmax variant `DFlashPagedBlockAttentionWarpKernel` (mirrors the shipped `AttentionWarpKernel`: one warp per (block-query,head), `__shfl_xor` butterfly reduction, register accumulator, NO `__syncthreads` storm; SAME paged/block combined-index read + causal/SWA mask + GQA; default ON, `VT_DFLASH_ATTN_BLOCK=1` keeps the bit-identical D12/D13 block kernel for A/B). Draft attn 242.9 → 77.9 ms (3.1x); our-ON c1 28.60 → 29.32 tok/s (+2.5%). **FINAL same-session 3-rep A/B (8 prompts×256 tok, cold leg discarded): our-ON graphed 29.42/29.27/29.32 (med 29.32) vs vLLM-ON graphed 29.240/29.247/29.233 (med 29.240) — our WORST rep (29.27) > vLLM's BEST (29.247), NON-OVERLAPPING bands, 1.003× ⇒ the ≥vLLM speed gate is MET.** Correctness UNCHANGED (output is exact by spec-decode construction — the target verify is untouched, only which draft proposals are accepted can shift): e2e `test_qwen27_dflash_spec_decode` 27/27 with graph==eager BIT-IDENTICAL, acceptance 19/39/29/25 unchanged (draft accepted 1629 identical warp-vs-block across the whole A/B set), 2/4 STRICT (France@11, 17×23@12 unchanged); CUDA==CPU `test_ops_dflash_paged_block_attn` 795648/795648 (warp within the f32 1e-4 / bf16 3e-2 envelope) + compute-sanitizer 0. Inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean; `check-device-leakage` not increased. `benchmark_binding=true`. Block-diffusion drafting is now correctness-complete (ratified near-tie) AND at/above vLLM throughput — this was the roadmap's FINAL open speed item. Anchors: `src/vt/cuda/cuda_ops.cu` (`DFlashPagedBlockAttentionWarpKernel` + `UseDflashAttnBlockKernel`; the D12 block kernel retained as the `VT_DFLASH_ATTN_BLOCK=1` reference). | T1 | `vllm/v1/worker/gpu/spec_decode/dflash/speculator.py`; `vllm/model_executor/models/qwen3_dflash.py`; `vllm/model_executor/models/interfaces.py:1382` (aux value); `eagle3_utils.py:41-56` (+1 shift) | `include/vllm/model_executor/models/qwen3_5.h` (`Qwen3_5AuxTaps`, `ForwardDeviceMultiTap`); `qwen3_5_dense.h`; `model_registry.h` (`aux_tap`); `src/vllm/model_executor/models/qwen3_5.cpp` (`MaybeCaptureAuxTap`/`ValidateAuxTapLayerIds`/`ForwardDeviceMultiTap`); `qwen3_5_moe.cpp`+`qwen3_5_dense.cpp` (routing); D2/D3 `include/vllm/model_executor/models/qwen3_dflash.h` + `src/vllm/model_executor/models/qwen3_dflash{,_weights}.cpp`; D4 `include/vllm/v1/worker/gpu/spec_decode/dflash/speculator.h` + `src/vllm/v1/worker/gpu/spec_decode/dflash/speculator.cpp` (`DflashProposeBlock`/`SampleDflashBlockDrafts`); D5 `src/vllm/entrypoints/model_loader.cpp` (`LoadDflashDraft`/`DflashDraft`) + `include/vllm/entrypoints/model_loader.h`; D5 `src/vllm/v1/worker/gpu/runner.cpp` (`set_dflash_draft`/`propose_drafts_dflash`/aux-tap capture) + `include/vllm/v1/worker/gpu/runner.h`; `src/vllm/config/speculative.cpp` + `include/vllm/config/speculative.h` (`ResolveDflash` + `dflash`/`model` parse); D14 warp kernel [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu#L1433) | `tests/vllm/models/test_qwen27_paged_forward.cpp` (multi-tap 598); `tests/vt/test_ops_dflash_block_attn.cpp`; `tests/vllm/models/test_qwen3_dflash_forward.cpp`; `tests/vllm/v1/spec_decode/test_dflash_kvprep.cpp`; `tests/parity/test_qwen3_dflash_{draft,kvprep}_parity.cpp`; D4 `tests/vllm/v1/spec_decode/test_dflash_propose.cpp` (5/19, RED-first); D5 `tests/parity/test_qwen27_dflash_spec_decode.cpp` (e2e 27/27, 2/4 strict + acceptance~vLLM); `scripts/spec/d{0,2,3}_dflash_*.py`; `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`; D6 `scripts/spec/vllm_dflash_timing.py` (vLLM-DFlash c1 timing); D7 device-resident `src/vllm/model_executor/models/qwen3_dflash.cpp` (`PrecomputeContextKVDevice` + `ForwardBlockLogitsWithContext` via `vt::IndexCopy`/`IndexSelect`); D9 persistent paged draft-KV `qwen3_dflash.{h,cpp}` (`AppendContextKVHost`/`ForwardBlockLogitsWithPrecomputedKV`/`ForwardWithCtxKVDev`/`PrecomputedContextKV`) + `runner.{h,cpp}` (`dflash_kv_store_`/`propose_drafts_dflash`) + `tests/vllm/v1/spec_decode/test_dflash_propose.cpp` (2 D9 bit-identity cases); D12 A-wire `runner.{h,cpp}` (device store as production path) + D12 Part B `include/vt/ops.h`/`src/vt/ops.cpp`/`src/vt/cpu/cpu_ops.cpp`/`src/vt/cuda/cuda_ops.cu` (`kDFlashPagedBlockAttention`) + `tests/vt/test_ops_dflash_paged_block_attn.cpp` (CPU==CUDA + cross-check, 795648/795648 + sanitizer-0); D13 `src/vllm/model_executor/models/qwen3_dflash.cpp` (fixed-capacity paged `DflashDeviceKVStore` + `ForwardPagedBody` + the per-request draft-step CUDA graph in `ForwardBlockLogitsWithDeviceKV`); D14 [test_ops_dflash_paged_block_attn](../tests/vt/test_ops_dflash_paged_block_attn.cpp#L79) + [ledger](parity-ledger.md#L738) | [dflash-spec-decode.md](specs/dflash-spec-decode.md) | `DONE` | `489a7544` |
 | `SPEC-DSPARK` | DSpark semi-autoregressive block drafter for DeepSeek-V4 and Qwen3, including native and Speculators checkpoint layouts, anchor-vs-bonus-token semantics, reduced/heterogeneous vocabulary mapping, sequential Markov sampling, noncausal draft attention, rejection/metrics and full-CUDA-graph compatibility; user-promoted scope at the v0.25.0 audit | T1 | `vllm/v1/worker/gpu/spec_decode/dspark/speculator.py:3-99`; `vllm/model_executor/models/qwen3_dspark.py:3-172`; `vllm/models/deepseek_v4/nvidia/dspark.py:56-457`; `tests/v1/e2e/spec_decode/test_spec_decode.py:1415-1469` @ `702f481` | - | - | `planned: specs/dspark-spec-decode.md` | `INVENTORIED` | - |
 | `SPEC-TLI` | Tokenizer-agnostic speculative decoding across heterogeneous draft/target vocabularies: shared-token mapping, target↔draft ID translation, constrained draft logits and greedy-only validation | T1 | `vllm/config/speculative.py:145-149,1173-1203`; `vllm/v1/spec_decode/vocab_mapping.py:68-160`; `vllm/v1/spec_decode/draft_model.py:34-58`; `vllm/v1/spec_decode/llm_base_proposer.py:432-495,688-691,831-837`; `tests/v1/spec_decode/test_vocab_mapping.py:1-50` @ `702f481` | - | - | `planned: specs/tli-spec-decode.md` | `INVENTORIED` | - |
 | `SPEC-NGRAM` | Draft-FREE n-gram proposer. **DONE 2026-07-27 (`CLAIM-ROADMAP-D3`):** 1:1 port of `ngram_proposer.py` (KMP-LPS suffix-ngram matcher + batch propose) wired as a third `--speculative-config` method reusing the LANDED MTP/DFlash verify/reject/`take_draft_token_ids` loop (no draft model / hidden tap / draft KV; GDN spec verify reused via `MakeQwen3_5KVCacheSpec(num_spec>0)`). 27B gate 5/5 STRICT our-ngram-ON == vLLM-ngram-ON + 180/180 drafts accepted; unit 19/19; spec-OFF byte-identical; host-side, no new kernel, `-Werror` clean | T2 | `vllm/v1/spec_decode/ngram_proposer.py:184-276,128-180`; `vllm/config/speculative.py:734-762,1224-1234`; `tests/v1/spec_decode/test_ngram.py` @ `555967922` | `src/vllm/v1/spec_decode/ngram_proposer.{h,cpp}`; `include/vllm/config/speculative.h` (`ResolveNgram`/`use_ngram`); `src/vllm/config/speculative.cpp`; `src/vllm/entrypoints/model_loader.cpp` (`ResolveSpecConfig`); `src/vllm/v1/worker/gpu/runner.cpp` (`propose_drafts_ngram`) | `tests/vllm/v1/spec_decode/test_ngram_proposer.cpp` (19/19); `tests/parity/test_qwen27_ngram_spec_decode.cpp` (5/5 STRICT, 180/180 accepted, dgx); golden `tests/parity/goldens/ngram_27b/ngram_27b_spec_on.json` + `scripts/spec/ngram_27b_golden.py`; ledger [parity-ledger.md](parity-ledger.md) 2026-07-27 | [specs/spec-decode-breadth-d3.md](specs/spec-decode-breadth-d3.md) | `ACTIVE` | `CLAIM-ROADMAP-D3` |
diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md
index e6c1997c..ce8fcef9 100644
--- a/.agents/kernel-matrix.md
+++ b/.agents/kernel-matrix.md
@@ -112,8 +112,8 @@ at 178.40 us/call. Evidence:
 | `KERNEL-ATTN-FLASHINFER-TRTLLM` | FlashInfer paged attention, XQA, and TensorRT-LLM generation kernels | availability/dispatch `vllm/utils/flashinfer.py:206-342,373-511`; SM90 decode and SM100+ prefill/decode gates | - | - | [inventory](specs/kernel-family-inventory.md) | `INVENTORIED` | - |
 | `KERNEL-ATTN-TRITON-FLEX-HPC` | Triton, FlexAttention, TurboQuant, and HPC attention families | registry `registry.py:34-120`; Triton contract `triton_attn.py:250-383`; HPC gate `hpc_attn.py:235-247` | - | - | [inventory](specs/kernel-family-inventory.md) | `INVENTORIED` | - |
 | `KERNEL-ATTN-MLA-SPARSE` | MLA and sparse attention: CUTLASS, FlashMLA, FlashInfer, FA, Triton, MSA **W6: the MLA attention BLOCK + WEIGHT ABSORPTION — the layer that COMPOSES W3+W4+W5** — [mla_attention.h](../include/vllm/model_executor/models/mla_attention.h) + [mla_attention.cpp](../src/vllm/model_executor/layers/attention/mla_attention.cpp) <- `mla.py:119-181` (`MultiHeadLatentAttentionWrapper.forward`) over `mla_attention.py:553-620` (the cache-update-BEFORE-attention order), `:624-874` (`forward_impl`: the dispatch + the absorbed decode) and `:2344-2425` (`forward_mha`); `AbsorbKvBProjBf16` <- `:875-962 process_weights_after_loading` (split `:892-900`, permutes `:959-962`); `MakeMlaUpProjectFn` <- `:2141-2170` (the `kv_b_proj` callback W5 left open); `BuildDeepseekRopeCosSinCache` <- `deepseek_scaling_rope.py:76-118` over `rotary_embedding/common.py:34-70`; `MlaAttentionScale` <- `deepseek_v2.py:995,1067-1075` (the mscale^2 correction, kept SEPARATE from the rope's own rotation mscale). **Absorption is a LOAD-TIME weight transform plus TWO batched GEMMs, not a fused kernel**, so it needed only two new general primitives: **`vt::BatchedMatmul`** <- `torch.bmm` at `mla_attention.py:789` (q-side W_UK fold) and `:1034` (`_v_up_proj`), CUDA impl = cuBLASLt STRIDED-BATCHED [cuda_matmul.cu](../src/vt/cuda/cuda_matmul.cu) (the cuBLASLt form of the cuBLAS `gemmStridedBatchedEx` torch.bmm resolves to; the only upstream alternatives are ROCm-only aiter fp8/fp4 bmm branches) + CPU ref [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp), stride-driven because BOTH call sites pass `.transpose(0,1)` views; and **`vt::ConcatMlaNopeRope`** <- `ConcatMLAQKernel` (`csrc/libtorch_stable/concat_mla_q.cuh`) + wrapper `cache_kernels.cu:1555-1600`, GENERALIZED to arbitrary nope/rope widths and a head-BROADCAST rope operand so one op also serves `_concat_k_nope_k_pe` (`:2063-2092`) — CUDA [cuda_mla_attn.cu](../src/vt/cuda/cuda_mla_attn.cu), CPU [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp). Two ADDITIVE relaxations of existing ops, integer-identical for contiguous tensors: `vt::RopeFromCache` stride-driven on q/k (DeepSeek rotates the TRAILING 64-dim slice and its `k_pe` is a column block of the fused kv_a projection) and `vt::MatmulBT` accepting a row-strided ACTIVATION (`kv_b_proj` applied to a 512-column slice of the 576-wide workspace, `:2160`) | CUDA priority `vllm/platforms/cuda.py:84-176` (`_get_backend_priorities`, both branches); MLA classes `vllm/v1/attention/backends/mla/*.py`; MLA prefill selector `mla/prefill/selector.py:47-76`; capability filter `vllm/v1/attention/backend.py:307-360`; CUTLASS build `CMakeLists.txt:1037-1061` **W6** [test_mla_attention_block.cpp](../tests/vllm/model_executor/layers/attention/test_mla_attention_block.cpp) **10/10 cases / 2,372,644 assertions** and [test_ops_mla_absorb.cpp](../tests/vt/test_ops_mla_absorb.cpp) **9/9 / 1,644,807 assertions** on dgx sm_121 — ports of `tests/kernels/test_concat_mla_q.py` (BOTH arms incl. the NON-CONTIGUOUS transposed-nope case, compared bit-exactly since a concat is a pure copy), the MLA-geometry sweep of `tests/v1/attention/test_mla_backends.py`, and the two-pass-oracle discipline of `tests/kernels/attention/test_mla_decode_cpu.py`. **THE ABSORBED-vs-UNABSORBED EQUIVALENCE IS PROVEN NUMERICALLY, THREE WAYS:** an INDEPENDENT double-precision block oracle computing the attention BOTH ways agrees to **< 1e-11** (the identity itself, at both query branches); our absorbed decode reproduces the UNABSORBED oracle to **< 2e-4** (f32); and the SAME batch driven once through the ABSORBED MQA decode kernel and once through the UNABSORBED materialized-MHA prefill path agrees to **< 3e-4** (CPU f32) / **< 4e-2** (CUDA bf16) — two code paths sharing nothing but the weights. Real geometry throughout (V2-Lite 512/128/64/128/16-head, plus V3's 7168 / 128-head / `q_lora_rank=1536` for the lora branch, which has NO e2e coverage and says so). Decode-only / prefill-only-no-context / chunked-prefill-with-context / MIXED (decode packed FIRST) all gated; NaN-poisoned outputs; run-to-run BIT-exact; CUDA cases proven to EXECUTE (124,941 + 290,835 assertions when run alone). memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors** (the last requires `--num-cuda-barriers 65536`: the default table OVERFLOWS on a binary driving this many kernel families and the tool then emits a bogus `unspecified launch failure`). Clean CUDA build 0 warn/0 err; regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 138/138, Qwen3-dense 664/664, OPT 36/36). No speed number — W9 owns tuning | **SELECTION (W2) + the DECODE KERNEL (W4) + the PREFILL PATH and CHUNKED-CONTEXT LOOP (W5). What is still absent is the MLA attention BLOCK and MODEL (W6/W7).** Priority TABLE [cuda_attn_priority.h:49](../include/vllm/platforms/cuda_attn_priority.h#L49) (both branches, one row per upstream arch arm) + lookup [cuda_attn_priority.h:86](../include/vllm/platforms/cuda_attn_priority.h#L86); `is_mla()`/`is_sparse()` filter [registry.cpp:63](../src/vllm/v1/attention/registry.cpp#L63); `TritonMLABackend` NAME + 3-D `get_kv_cache_shape` [backend.h:271](../include/vllm/v1/attention/backend.h#L271), [backend.cpp:83](../src/vllm/v1/attention/backend.cpp#L83), registration [backend.cpp:108](../src/vllm/v1/attention/backend.cpp#L108) — **W4: `vt::MlaDecodeAttention`** — `OpId::kMlaDecodeAttention` + args/validation [ops.h](../include/vt/ops.h), [ops.cpp](../src/vt/ops.cpp); CPU single-pass REFERENCE [cpu_mla_attn.cpp](../src/vt/cpu/cpu_mla_attn.cpp) (numerics from `csrc/cpu/mla_decode.cpp`); CUDA two-stage split-KV [cuda_mla_attn.cu](../src/vt/cuda/cuda_mla_attn.cu) — `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, IS_MLA `v = tl.trans(k)` branch `:424-431`), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace via the house grow-only per-stream scratch (upstream's `_reserve_attn_logits_workspace` `:57-78`). Deterministic by construction: fixed ASCENDING split merge, NO atomicAdd. `TritonMLABackend::get_impl_cls()` now returns a real `TritonMLAImpl` [backend.h](../include/vllm/v1/attention/backend.h), [backend.cpp](../src/vllm/v1/attention/backend.cpp); PREFILL remains W5 and `TritonMLAImpl::forward` refuses a prefill-shaped batch by name. **W5: `vt::MlaPrefillAttention` + `vt::GatherMlaCache` + `vt::MergeAttnStates` + the chunked-context driver** — `vt::MlaPrefillAttention` [cuda_mla_prefill.cu](../src/vt/cuda/cuda_mla_prefill.cu) / CPU ref [cpu_mla_prefill.cpp](../src/vt/cpu/cpu_mla_prefill.cpp) <- `mla/prefill/flash_attn.py:153-248` `FlashAttnPrefillBackend` (the ONLY MLA prefill backend reachable on sm_121 per `mla/prefill/selector.py:66-76`, and it HARD-RAISES with no fallback at `:191-194`), running over the vendored FA-2 through the NEW launcher entry `LaunchMlaPrefillFA2Bf16` [cuda_flash_attn_fa2.cu](../src/vt/cuda/cuda_flash_attn_fa2.cu) plus two new explicit instantiations of the UNCHANGED generic template (`flash_fwd_split_hdim192_bf16{,_causal}_sm80.cu`). V is ZERO-PADDED 128->192 and the output sliced back, exactly as upstream's `requires_v_padding` path does (`flash_attn.py:88-99,164-168,196-197`) — which is WHY the asymmetric QK 192 / V 128 pair needs no asymmetric kernel. `vt::GatherMlaCache` <- `csrc/libtorch_stable/cache_kernels.cu:992-1064`; `vt::MergeAttnStates` <- `csrc/libtorch_stable/attention/merge_attn_states.cu:18-192` (BOTH `-inf` edge cases ported verbatim). The workspace-bounded loop is [mla_chunked_context.h](../include/vllm/model_executor/layers/attention/mla_chunked_context.h) <- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`. **The paged launcher `LaunchPrefillFA2Bf16` that every non-MLA prefill calls is textually UNTOUCHED** (211 insertions / 0 deletions in that TU; 2 new vendored files) | [test_attn_backend_registry.cpp:146](../tests/vllm/v1/attention/test_attn_backend_registry.cpp#L146) (GB10 MLA list), [:203](../tests/vllm/v1/attention/test_attn_backend_registry.cpp#L203) (`use_mla=true` -> `TRITON_MLA`, matching the W0 oracle observation), [:230](../tests/vllm/v1/attention/test_attn_backend_registry.cpp#L230) (the DSA seam, proven both directions with a stand-in sparse backend) — ports of `test_attention_backends_selection.py` (MLA cases), `test_mla_prefill_selector.py`, `test_mla_prefill_registry.py`; **W4** [test_ops_mla_attn.cpp](../tests/vt/test_ops_mla_attn.cpp) — port of `tests/kernels/attention/test_mla_decode_cpu.py` (`ref_mla` as a TWO-PASS oracle, its bs=4/mean_seq_len=256/h_q=16/d=576/dv=512/block=16 parametrization, BOTH varlen arms, and its NaN-padding out-of-bounds detector) plus the `test_mla_backends.py` shape sweep: ragged, multi-block, single-block/single-token, EVERY num_kv_splits in {1,2,3,4,5,8,16,17,64,300,512} (incl. splits > seq_len, the empty-split path both stages must skip), 128-head DeepSeek-V3 geometry, head counts 1/3/17 that do not fill a BLOCK_H tile, a 288/256 block-32 non-V2-Lite geometry, bf16 + f32, and run-to-run BIT-exactness over 5 runs. Gated on dgx/sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build 0 warn/0 err; regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6). NO speed number yet — decode perf is W9. **W5** [test_ops_mla_prefill.cpp](../tests/vt/test_ops_mla_prefill.cpp) **4/4 cases / 2,377,052 assertions** and [test_ops_mla_chunked_context.cpp](../tests/vt/test_ops_mla_chunked_context.cpp) **5/5 / 306,037 assertions** on dgx sm_121 — ports of `tests/v1/attention/test_mla_backends.py` and `tests/v1/attention/test_mla_prefill_quant_output.py` (its fp8 arms NOT ported: they need device-capability family 100, unreachable on sm_121 — recorded, not dropped). REAL V2-Lite prefill geometry (QK 192 / V 128 / latent 576, block 16, mscale^2 scale) against an INDEPENDENT double-precision TWO-PASS oracle, plus — for the chunked loop — a SINGLE-SHOT whole-sequence oracle that never chunks: exact / +1 / -1 chunk boundaries, a request with NO context, a chunk in which a request contributes ZERO keys, ragged multi-chunk, 128-head V3, single-token queries, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs, run-to-run BIT-exact over 5 runs. memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors** on both binaries; clean CUDA build 0 warn/0 err; regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 138/138, Qwen3-dense 664/664, OPT 36/36). Prefill perf is W9 | [MLA campaign spike](specs/mla-deepseek-campaign.md) | `PARTIAL` | `CLAIM-MLA-DEEPSEEK` |
-| `KERNEL-ATTN-DFLASH-BLOCK` | **DFlash in-block attention — the project's FIRST non-causal / bidirectional attention primitive** (SPEC-DFLASH D2, DF-DRAFT-MODEL). Per-request uniform (1+k) query block attends within its own block: FULL-attention layers BIDIRECTIONAL (`causal=false`, no mask), SWA layers causal-within-window. f32 online softmax, GQA broadcast. A SEPARATE `vt::` op from the causal `kAttention`/`kPagedAttention` so every other model stays byte-identical | `vllm/model_executor/models/qwen3_dflash.py:86-146` (`_resolve_layer_attention`: full layers default non-causal, SWA causal) + `:149-263` (`DFlashQwen3Attention`); flashinfer non-causal path (vllm#48167 Blackwell non-causal attn, in-pin) | `OpId::kDFlashBlockAttention` + `DFlashBlockAttentionArgs` + decl [ops.h:1713](../include/vt/ops.h#L1713) + wrapper/validation [ops.cpp:2069](../src/vt/ops.cpp#L2069); CPU REFERENCE `DFlashBlockAttentionKernel` [cpu_ops.cpp:1843](../src/vt/cpu/cpu_ops.cpp#L1843) (three-pass block-local softmax, the authoritative impl); CUDA `DFlashBlockAttentionKernelCuda` [cuda_ops.cu:1300](../src/vt/cuda/cuda_ops.cu#L1300) mirroring the causal `AttentionKernel` block-reduction recurrence with per-block bounds + the bidirectional/window mask; the draft model that consumes it [qwen3_dflash.cpp:52](../src/vllm/model_executor/models/qwen3_dflash.cpp#L52) | **CPU GATE GREEN** [test_ops_dflash_block_attn.cpp:79](../tests/vt/test_ops_dflash_block_attn.cpp#L79) 5 cases / 12 assertions — hand-checked non-causal (query 0 sees the future key), the RED causal-vs-non-causal separation (the mask is load-bearing), per-request cu_seqlens block isolation, SWA window bound, GQA; model forward [test_qwen3_dflash_forward.cpp:116](../tests/vllm/models/test_qwen3_dflash_forward.cpp#L116) 5 cases / 95 assertions (RED full-layer-causal-flip); existing causal `test_ops_attention` 9/9·23 UNCHANGED. **GPU GATE GREEN on dgx (2026-07-26, GB10 sm_121a):** CUDA `-Werror=all-warnings` build clean (kernel compiles as-written, no change); CUDA==CPU parity [test_ops_dflash_block_attn CUDA case](../tests/vt/test_ops_dflash_block_attn.cpp#L153) 198412/198412 within the 1e-4 f32-softmax envelope over all 5 corners; `compute-sanitizer --tool memcheck` 0 errors; consumed by the draft-forward parity gate ([test_qwen3_dflash_draft_parity](../tests/parity/test_qwen3_dflash_draft_parity.cpp), fc rel-L2 0.46% / hidden ≤1.3% vs the real vLLM draft). **DONE 2026-07-27 with the DFlash block (`CLAIM-DFLASH-D14`):** the D2 non-causal in-block primitive is the CPU/materialized reference the D12+ paged/warp kernels are gated against; closure [ledger](parity-ledger.md#L722). | [DFlash spec](specs/dflash-spec-decode.md) §1.3/§6 D2 | `DONE` | `164453a2` |
-| `KERNEL-ATTN-DFLASH-PAGED-BLOCK` | **DFlash PAGED in-block attention — the CAPTURE-SAFE form of `KERNEL-ATTN-DFLASH-BLOCK`** (SPEC-DFLASH D12 Part B, the CUDA-graph draft-attention primitive). The (1+k) block queries attend over `[PAGED context ; their own (1+k) block]`: the growing context enters as DATA (paged K/V cache `[pages,block_size,Hkv,D]` + per-request `seq_lens` + `block_table`, mirroring `PagedAttentionKernel`) instead of a variable-size materialized combined buffer, so the launch grid is STATIC over the fixed `Nq=(1+k)*num_reqs` rows and EVERY metadata input is a persistent DEVICE tensor read in place — NO `cudaMallocAsync`/`cudaMemcpyAsync` of a function-local host `cu_seqlens` (the [[cudagraph-capture-bakes-stack-addresses]] UAF class the eager `LaunchDFlashBlockAttention` had). Same f32 online softmax + D2 in-block mask over the COMBINED index; bit-identical to `DFlashBlockAttention` over the materialized `[context;block]` buffer | vLLM full CG `dflash/cudagraph.py` + `speculator.py:411-458` + `precompute_and_store_context_kv` (`qwen3_dflash.py:548-619`) @ `555967922`; paged read mirrors our `PagedAttentionKernel` [cuda_paged_attn.cu:184](../src/vt/cuda/cuda_paged_attn.cu#L184) | `OpId::kDFlashPagedBlockAttention` + `DFlashPagedBlockAttentionArgs` + decl [ops.h](../include/vt/ops.h) + wrapper/validation [ops.cpp](../src/vt/ops.cpp); CPU REFERENCE `DFlashPagedBlockAttentionKernel` [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp); CUDA `DFlashPagedBlockAttentionKernelCuda` [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu#L1452) (static grid, persistent device metadata) + D14 WARP variant [DFlashPagedBlockAttentionWarpKernel](../src/vt/cuda/cuda_ops.cu#L1433) | **GPU GATE GREEN on dgx (2026-07-27, GB10 sm_121a):** CUDA `-Werror` clean (0 warnings); [test_ops_dflash_paged_block_attn.cpp](../tests/vt/test_ops_dflash_paged_block_attn.cpp#L79) cross-checks CPU-paged == materialized `DFlashBlockAttention` across 6 corners (non-causal, causal-SWA, block isolation, GQA, multi-page, zero-context) + CUDA==CPU (f32+bf16) = **795648/795648 assertions**; `compute-sanitizer --tool memcheck` **0 errors**. **D13 (2026-07-27, `CLAIM-DFLASH-D13`): WIRED INTO PRODUCTION** — the single-request DFlash draft forward (`ForwardPagedBody`, `qwen3_dflash.cpp`) now runs the (1+k) block through this kernel reading a fixed-capacity paged `DflashDeviceKVStore`, and the whole draft step is captured into a per-request CUDA graph + replayed (the growing context enters only via the in-place `seq_lens`). Capture-correctness PROVEN: `test_qwen27_dflash_spec_decode` 27/27 with the graph BIT-IDENTICAL to eager (same tokens + acceptance 19/39/29/25); c1 throughput NEAR-PARITY with vLLM-DFlash-ON (ours 0.978×, ~2% below the tight 3-rep band; gap closed 0.917×→0.978× via the paged read, the CG is perf-neutral) — the kernel is landed + wired + gated; STAYS `ACTIVE` with the engine feature (the ~2% ≥vLLM residual is per-step compute for an nsys). **D14 (2026-07-27, `CLAIM-DFLASH-D14`): the residual WAS this kernel → WARP-scoped variant added → SPEED GATE MET, `DONE`.** An nsys (`--cuda-graph-trace=node`) attributed the D13 ~2% residual to THIS kernel: `DFlashPagedBlockAttentionKernel` = 242.9 ms = 1.8% of the graphed step's GPU time, median ~460 us/call (grid `(nq=17,hq=32)` × kBlock=256 threads looping SERIALLY over C~500-640 keys with a 256-wide shared-mem tree reduction + 2 `__syncthreads` PER key — the latency/sync storm the ViT tower fixed with `AttentionDenseFast`), vs vLLM's fused flash draft-attn ~0.15%. Added `DFlashPagedBlockAttentionWarpKernel` ([cuda_ops.cu](../src/vt/cuda/cuda_ops.cu)): ONE WARP per (block-query,head), `__shfl_xor` butterfly head_dim reduction, register accumulator, NO `__syncthreads`; SAME paged/block combined-index read + causal/SWA mask + GQA (copied verbatim from the block kernel), mirroring the shipped `AttentionWarpKernel`. Default ON; `VT_DFLASH_ATTN_BLOCK=1` keeps the bit-identical D12/D13 block kernel. Draft attn **242.9 → 77.9 ms (3.1×)**; our-ON c1 **28.60 → 29.32 tok/s**; FINAL 3-rep A/B our-ON 29.32 ≥ vLLM-ON 29.240 (non-overlapping bands, 1.003×) ⇒ **≥vLLM MET**. Not bit-identical to the block kernel but same f32-online-softmax math within envelope; CUDA==CPU `test_ops_dflash_paged_block_attn` **795648/795648** (f32 1e-4/bf16 3e-2) + **compute-sanitizer 0**; e2e 27/27 graph==eager, acceptance 19/39/29/25 unchanged (1629 accepted identical warp-vs-block); SACRED 235/235 + MTP 9/9 inert; `-Werror` clean; closure [ledger](parity-ledger.md#L738) | [DFlash spec](specs/dflash-spec-decode.md) §0 D12/D13/D14 | `DONE` | `164453a2` |
+| `KERNEL-ATTN-DFLASH-BLOCK` | **DFlash in-block attention — the project's FIRST non-causal / bidirectional attention primitive** (SPEC-DFLASH D2, DF-DRAFT-MODEL). Per-request uniform (1+k) query block attends within its own block: FULL-attention layers BIDIRECTIONAL (`causal=false`, no mask), SWA layers causal-within-window. f32 online softmax, GQA broadcast. A SEPARATE `vt::` op from the causal `kAttention`/`kPagedAttention` so every other model stays byte-identical | `vllm/model_executor/models/qwen3_dflash.py:86-146` (`_resolve_layer_attention`: full layers default non-causal, SWA causal) + `:149-263` (`DFlashQwen3Attention`); flashinfer non-causal path (vllm#48167 Blackwell non-causal attn, in-pin) | `OpId::kDFlashBlockAttention` + `DFlashBlockAttentionArgs` + decl [ops.h:1713](../include/vt/ops.h#L1713) + wrapper/validation [ops.cpp:2069](../src/vt/ops.cpp#L2069); CPU REFERENCE `DFlashBlockAttentionKernel` [cpu_ops.cpp:1843](../src/vt/cpu/cpu_ops.cpp#L1843) (three-pass block-local softmax, the authoritative impl); CUDA `DFlashBlockAttentionKernelCuda` [cuda_ops.cu:1300](../src/vt/cuda/cuda_ops.cu#L1300) mirroring the causal `AttentionKernel` block-reduction recurrence with per-block bounds + the bidirectional/window mask; the draft model that consumes it [qwen3_dflash.cpp:52](../src/vllm/model_executor/models/qwen3_dflash.cpp#L52) | **CPU GATE GREEN** [test_ops_dflash_block_attn.cpp:79](../tests/vt/test_ops_dflash_block_attn.cpp#L79) 5 cases / 12 assertions — hand-checked non-causal (query 0 sees the future key), the RED causal-vs-non-causal separation (the mask is load-bearing), per-request cu_seqlens block isolation, SWA window bound, GQA; model forward [test_qwen3_dflash_forward.cpp:116](../tests/vllm/models/test_qwen3_dflash_forward.cpp#L116) 5 cases / 95 assertions (RED full-layer-causal-flip); existing causal `test_ops_attention` 9/9·23 UNCHANGED. **GPU GATE GREEN on dgx (2026-07-26, GB10 sm_121a):** CUDA `-Werror=all-warnings` build clean (kernel compiles as-written, no change); CUDA==CPU parity [test_ops_dflash_block_attn CUDA case](../tests/vt/test_ops_dflash_block_attn.cpp#L153) 198412/198412 within the 1e-4 f32-softmax envelope over all 5 corners; `compute-sanitizer --tool memcheck` 0 errors; consumed by the draft-forward parity gate ([test_qwen3_dflash_draft_parity](../tests/parity/test_qwen3_dflash_draft_parity.cpp), fc rel-L2 0.46% / hidden ≤1.3% vs the real vLLM draft). **DONE 2026-07-27 with the DFlash block (`CLAIM-DFLASH-D14`):** the D2 non-causal in-block primitive is the CPU/materialized reference the D12+ paged/warp kernels are gated against; closure [ledger](parity-ledger.md#L722). | [DFlash spec](specs/dflash-spec-decode.md) §1.3/§6 D2 | `DONE` | `489a7544` |
+| `KERNEL-ATTN-DFLASH-PAGED-BLOCK` | **DFlash PAGED in-block attention — the CAPTURE-SAFE form of `KERNEL-ATTN-DFLASH-BLOCK`** (SPEC-DFLASH D12 Part B, the CUDA-graph draft-attention primitive). The (1+k) block queries attend over `[PAGED context ; their own (1+k) block]`: the growing context enters as DATA (paged K/V cache `[pages,block_size,Hkv,D]` + per-request `seq_lens` + `block_table`, mirroring `PagedAttentionKernel`) instead of a variable-size materialized combined buffer, so the launch grid is STATIC over the fixed `Nq=(1+k)*num_reqs` rows and EVERY metadata input is a persistent DEVICE tensor read in place — NO `cudaMallocAsync`/`cudaMemcpyAsync` of a function-local host `cu_seqlens` (the [[cudagraph-capture-bakes-stack-addresses]] UAF class the eager `LaunchDFlashBlockAttention` had). Same f32 online softmax + D2 in-block mask over the COMBINED index; bit-identical to `DFlashBlockAttention` over the materialized `[context;block]` buffer | vLLM full CG `dflash/cudagraph.py` + `speculator.py:411-458` + `precompute_and_store_context_kv` (`qwen3_dflash.py:548-619`) @ `555967922`; paged read mirrors our `PagedAttentionKernel` [cuda_paged_attn.cu:184](../src/vt/cuda/cuda_paged_attn.cu#L184) | `OpId::kDFlashPagedBlockAttention` + `DFlashPagedBlockAttentionArgs` + decl [ops.h](../include/vt/ops.h) + wrapper/validation [ops.cpp](../src/vt/ops.cpp); CPU REFERENCE `DFlashPagedBlockAttentionKernel` [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp); CUDA `DFlashPagedBlockAttentionKernelCuda` [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu#L1452) (static grid, persistent device metadata) + D14 WARP variant [DFlashPagedBlockAttentionWarpKernel](../src/vt/cuda/cuda_ops.cu#L1433) | **GPU GATE GREEN on dgx (2026-07-27, GB10 sm_121a):** CUDA `-Werror` clean (0 warnings); [test_ops_dflash_paged_block_attn.cpp](../tests/vt/test_ops_dflash_paged_block_attn.cpp#L79) cross-checks CPU-paged == materialized `DFlashBlockAttention` across 6 corners (non-causal, causal-SWA, block isolation, GQA, multi-page, zero-context) + CUDA==CPU (f32+bf16) = **795648/795648 assertions**; `compute-sanitizer --tool memcheck` **0 errors**. **D13 (2026-07-27, `CLAIM-DFLASH-D13`): WIRED INTO PRODUCTION** — the single-request DFlash draft forward (`ForwardPagedBody`, `qwen3_dflash.cpp`) now runs the (1+k) block through this kernel reading a fixed-capacity paged `DflashDeviceKVStore`, and the whole draft step is captured into a per-request CUDA graph + replayed (the growing context enters only via the in-place `seq_lens`). Capture-correctness PROVEN: `test_qwen27_dflash_spec_decode` 27/27 with the graph BIT-IDENTICAL to eager (same tokens + acceptance 19/39/29/25); c1 throughput NEAR-PARITY with vLLM-DFlash-ON (ours 0.978×, ~2% below the tight 3-rep band; gap closed 0.917×→0.978× via the paged read, the CG is perf-neutral) — the kernel is landed + wired + gated; STAYS `ACTIVE` with the engine feature (the ~2% ≥vLLM residual is per-step compute for an nsys). **D14 (2026-07-27, `CLAIM-DFLASH-D14`): the residual WAS this kernel → WARP-scoped variant added → SPEED GATE MET, `DONE`.** An nsys (`--cuda-graph-trace=node`) attributed the D13 ~2% residual to THIS kernel: `DFlashPagedBlockAttentionKernel` = 242.9 ms = 1.8% of the graphed step's GPU time, median ~460 us/call (grid `(nq=17,hq=32)` × kBlock=256 threads looping SERIALLY over C~500-640 keys with a 256-wide shared-mem tree reduction + 2 `__syncthreads` PER key — the latency/sync storm the ViT tower fixed with `AttentionDenseFast`), vs vLLM's fused flash draft-attn ~0.15%. Added `DFlashPagedBlockAttentionWarpKernel` ([cuda_ops.cu](../src/vt/cuda/cuda_ops.cu)): ONE WARP per (block-query,head), `__shfl_xor` butterfly head_dim reduction, register accumulator, NO `__syncthreads`; SAME paged/block combined-index read + causal/SWA mask + GQA (copied verbatim from the block kernel), mirroring the shipped `AttentionWarpKernel`. Default ON; `VT_DFLASH_ATTN_BLOCK=1` keeps the bit-identical D12/D13 block kernel. Draft attn **242.9 → 77.9 ms (3.1×)**; our-ON c1 **28.60 → 29.32 tok/s**; FINAL 3-rep A/B our-ON 29.32 ≥ vLLM-ON 29.240 (non-overlapping bands, 1.003×) ⇒ **≥vLLM MET**. Not bit-identical to the block kernel but same f32-online-softmax math within envelope; CUDA==CPU `test_ops_dflash_paged_block_attn` **795648/795648** (f32 1e-4/bf16 3e-2) + **compute-sanitizer 0**; e2e 27/27 graph==eager, acceptance 19/39/29/25 unchanged (1629 accepted identical warp-vs-block); SACRED 235/235 + MTP 9/9 inert; `-Werror` clean; closure [ledger](parity-ledger.md#L738) | [DFlash spec](specs/dflash-spec-decode.md) §0 D12/D13/D14 | `DONE` | `489a7544` |
 | `KERNEL-ATTN-DSA-SPARSE-INDEX` | **DeepSeek-V4 DSA "Lightning Indexer" sparse-attention SELECTION — the project's FIRST sparse candidate-selection primitive** (DeepSeek-V4-Flash W3). Two ops: (1) the weighted-MQA INDEXER LOGIT `logit[t,s] = Σ_h w[t,h]·ReLU(q[t,h]·k[s])` over the causal candidate window (the per-head **ReLU** is load-bearing — it is what makes the indexer a learned sparse SELECTOR, not a plain attention score), where `w[t,h] = weights_proj[t,h]·index_head_dim^-0.5·index_n_heads^-0.5`; and (2) the per-row **causal top-k** that keeps the `index_topk=512` highest-logit keys (short-context: every candidate, ascending; else top-k with -1 padding). Distinct from every dense/paged/MLA family, which score ALL keys — this one PICKS a sparse key subset the downstream MLA then attends over. W3 also lands the two 512-wide-MLA OUTPUT seams V2/V3 lack (per-head attention-**sink** softmax + **grouped output-LoRA** `wo_a` bmm→`wo_b`) as portable host references beside it | MQA logit `vllm/v1/attention/ops/triton_fp8_mqa_logits.py:120-156` (dot→×kv_scale→ReLU→×weights→Σheads); weight fold `vllm/model_executor/layers/sparse_attn_indexer.py:203-207`; top-k `sparse_attn_indexer.py:488-497` + short-context `vllm/models/deepseek_v4/attention.py:70-86,:813-831`; sinks `deepseek_v4/nvidia/flashinfer_sparse.py:777,:896`; grouped output-LoRA `deepseek_v4/nvidia/ops/o_proj.py:58-73` @ `555967922` | Portable host reference (device kernel is a W7 residual) [deepseek_v4_dsa.cpp](../src/vllm/model_executor/models/deepseek_v4_dsa.cpp) + [deepseek_v4_dsa.h](../include/vllm/model_executor/models/deepseek_v4_dsa.h): `DsaIndexerWeightFold` / `DsaIndexerLogits` / `DsaTopkSelect` / `SoftmaxWithSink` / `GroupedOutputLora` | **CPU UNIT GATE GREEN (2026-07-28, `-Wall -Werror -Wextra` 0-warn):** [test_deepseek_v4_dsa.cpp](../tests/vllm/models/test_deepseek_v4_dsa.cpp) **13/13 cases · 38 assertions** — hand-derived literal cases (the ReLU clip, the weight fold, short-context all-select, full top-k, tie→smaller-index, causal-window offset, sink probability mass, sink numerical stability, grouped-LoRA) + from-first-principles double-precision references on randomized shapes (indexer logits + grouped output-LoRA rel-L2 < 1e-6). Full-model gate is multi-Spark-blocked (156.7 GiB); MHC (W5) + sqrtsoftplus/hash MoE (W6) + device kernel + forward integration (W7) are named residuals | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W3 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W3` |
 | `KERNEL-ATTN-DSA-COMPRESSOR` | **DeepSeek-V4 DSA COMPRESSOR + fp8_ds_mla KV-cache state — the second half of the sparse-attention stack** (DeepSeek-V4-Flash W4). Where `KERNEL-ATTN-DSA-SPARSE-INDEX` SELECTS keys, this POOLS + QUANTIZES them into the compressed latent the MLA reads and defines how it is cached across steps. Three ops: **(1)** the softmax-weighted window POOL — at a compress boundary the compressor gathers `(1+overlap)·compress_ratio` KV-state rows and computes, PER head-dim column, `softmax(score, dim=0)·kv` (each channel pools the window with its OWN weights — the load-bearing nuance), then RMSNorm; **(2)** the fused save-time APE add `score_state = score + ape[position % compress_ratio]`; **(3)** the **fp8_ds_mla** KV-cache STATE layout — the 512-wide latent split into a 448-wide NoPE part quantized to FP8 e4m3 with per-64 **UE8M0** power-of-two block scales (exponent `= ceil(log2(absmax/448))`, byte `= exp+127`) and a 64-wide RoPE part stored bf16, at a **576-byte** token stride with a padded **7+1** scale region — plus the dequant READ (`nope = e4m3·2^(byte-127)`, `rope = bf16`) | pool+RMSNorm `vllm/models/deepseek_v4/common/ops/fused_compress_quant_cache.py:198-218`; save-time APE `common/ops/save_partial_states.py:92-101`; fp8_ds_mla store `fused_compress_quant_cache.py:220-297`; layout `deepseek_v4/compressor.py:307-309`; dequant READ cross-checked to SGLang `v0.5.15` `dsv4/dequant_k_cache.py:12-18,:122-136` @ `555967922` | Portable host reference (device kernel is a W7 residual) [deepseek_v4_compressor.cpp](../src/vllm/model_executor/models/deepseek_v4_compressor.cpp) + [deepseek_v4_compressor.h](../include/vllm/model_executor/models/deepseek_v4_compressor.h): `CompressorSaveScoreApe` / `CompressorPoolNorm` / `MakeFp8DsMlaLayout` / `Fp8DsMlaEncodeToken` / `Fp8DsMlaDecodeToken` | **CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, 0-warn on the new TUs):** [test_deepseek_v4_compressor.cpp](../tests/vllm/models/test_deepseek_v4_compressor.cpp) **12/12 cases · 164 assertions** — hand-derived literal cases (APE modulo wrap; per-column softmax pool proven load-bearing via the column-ratio-survives-RMSNorm case; window masking; V4 layout 448/64/576/7+1; all-ones→UE8M0 byte 119 exact round-trip; value-3→byte 120; bf16 rope verbatim) + from-first-principles double-precision references (pool+norm rel-L2 < 1e-6; independent UE8M0 scale-byte recompute; encode→decode round-trip < 0.05 fp8 granularity). RED-first PROVEN: perturbing the scale bias `+127→+126` fails 4 cases / 135 assertions; revert restores 12/12. Honest gate form: hand-case + structural review vs vLLM+SGLang `file:line` (fixed-config 167B not constructible at a tiny shape ⇒ NOT a dumped-oracle rel-L2). Full-model gate multi-Spark-blocked (156.7 GiB); MHC (W5) + sqrtsoftplus/hash MoE (W6) + the fused device kernel + forward integration (W7) are named residuals | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W4 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W4` |
 | `KERNEL-MHC-SINKHORN` | **DeepSeek-V4 Manifold/Markov Hyper-Connections (MHC) — the Sinkhorn-normalized hc_mult-stream residual manifold** (DeepSeek-V4-Flash W5, the hardest V4 brick). V4 replaces the plain `residual + RMSNorm` stream with a `[tokens, hc_mult=4, hidden]` MANIFOLD of parallel residual streams, mixed at every attn/ffn boundary by a **doubly-stochastic** matrix and collapsed by a learned head. Four ops: **(1)** the **`hc_sinkhorn_iters=20` Sinkhorn** normalization of the hc_mult×hc_mult mixing matrix — a row-softmax seed (`+eps`), a col-norm, then `(iters-1)×[row-norm, col-norm]` toward a doubly-stochastic matrix (the AXIS ALTERNATION and the ITERATION COUNT are load-bearing at non-converged counts — RED-first proven); **(2)** the mHC **pre** mix — flatten the streams, project through `hc_*_fn` with a FOLDED weight-free RMSNorm `rsqrt(sqrsum/(hc·H)+rms_eps)`, split into pre/post/comb gates (`pre=σ+hc_eps`, `post=σ·hc_post_alpha(2.0)`, `comb=Sinkhorn`), collapse to the single `layer_input`, and optionally FOLD the model's attn_norm/ffn_norm RMSNorm; **(3)** the mHC **post** mix — fold the block output back into the manifold via the comb matrix (`Σ_i comb[i,j]·res[i,h]`) + the post gate; **(4)** the **hc_head** collapse — weight-free RMSNorm → `hc_head_fn` → sigmoid gate → weighted stream sum → one hidden vector. **EAGER-REF FINDING: corrects the W0 "ZERO eager reference upstream" premise** — the pinned vLLM DOES ship an eager PyTorch reference (`mhc/torch.py` `mhc_pre_torch`/`mhc_post_torch`, `triton.py` head collapse); four upstream impls (torch.py, tilelang_kernels.py `_sinkhorn_fwd`, tilelang.py, SGLang mhc.py) agree byte-for-byte on the Sinkhorn | mHC pre/post + Sinkhorn `vllm/model_executor/kernels/mhc/torch.py:56-106` (byte-identical `tilelang_kernels.py:126-153` `_sinkhorn_fwd`, `tilelang.py` `mhc_pre_big_fuse_with_norm`); head collapse `triton.py:108-140` + `tilelang.py:720-748`; constants `hc_post_alpha=2.0`/`hc_pre_eps=hc_sinkhorn_eps=hc_eps` `vllm/models/deepseek_v4/nvidia/model.py:818-821,:886-894,:1023-1041`; cross-checked SGLang `v0.5.15` `python/sglang/srt/layers/mhc.py:110-126` @ `555967922` | Portable host reference (device kernel + `DeepseekV4Model::Forward` assembly are W7 residuals) [deepseek_v4_mhc.cpp](../src/vllm/model_executor/models/deepseek_v4_mhc.cpp) + [deepseek_v4_mhc.h](../include/vllm/model_executor/models/deepseek_v4_mhc.h): `MhcSinkhorn` / `MhcPre` / `MhcPost` / `HcHeadCollapse` | **CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, `-Wall -Werror -Wextra` 0-warn on the new TUs):** [test_deepseek_v4_mhc.cpp](../tests/vllm/models/test_deepseek_v4_mhc.cpp) **14/14 cases · 125 assertions** — hand-derived literal cases (all-zero Sinkhorn → uniform doubly-stochastic 1/hc; symmetric-2×2 fixed point `[[.75,.25],[.25,.75]]`; iteration-count load-bearing; MhcPre fn=0 gate midpoints; RMSNorm fold `[1,3]→[1,3]/√5`; MhcPost identity-comb + post-add; mix sums over the first comb index; hc_head fn=0 stream mean) + from-first-principles DOUBLE-PRECISION references (Sinkhorn/MhcPre/MhcPost/HcHead f32==f64 rel-L2 < 1e-5..1e-4; doubly-stochastic convergence to row/col sums=1). **RED-first PROVEN both levers:** perturb the Sinkhorn iteration count (`iters-1→iters-2`) fails 1 case/9 assertions AND swap a normalization axis fails 2 cases/12 assertions (caught by a dedicated SMALL-iteration-count gate, since at 20 iters the Sinkhorn has converged and ±1 is within tolerance); revert restores 14/14·125. Honest gate form: DERIVED-eager-reference + hand-case + structural review vs vLLM+SGLang `file:line` (fixed-config 167B not constructible at a tiny shape ⇒ NOT a dumped-oracle rel-L2). OPEN QUESTION: end-to-end bf16 residual/layer_input rounding between steps is a W7 device concern, not folded into these f32/f64 refs. Full-model gate multi-Spark-blocked (156.7 GiB); sqrtsoftplus/hash MoE (W6) + device kernel + forward assembly (W7) are named residuals | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W5 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W5` |
diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md
index b30f21d2..4c1d448c 100644
--- a/.agents/model-matrix.md
+++ b/.agents/model-matrix.md
@@ -474,7 +474,7 @@ Transformers compatibility is capability-driven and excluded from finite counts.
 | `MODEL-SPEC-llama-eagle-eagle-llama-for-causal-lm` | `EagleLlamaForCausalLM` | `registry.py:588`; `vllm/model_executor/models/llama_eagle.py::EagleLlamaForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; EAGLE/EAGLE3 | ☐ required | `INVENTORIED` | none | unassigned |
 | `MODEL-SPEC-llama4-eagle-eagle-llama4-for-causal-lm` | `EagleLlama4ForCausalLM` | `registry.py:589`; `vllm/model_executor/models/llama4_eagle.py::EagleLlama4ForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; EAGLE/EAGLE3 | ☐ required | `INVENTORIED` | none | unassigned |
 | `MODEL-SPEC-minicpm-eagle-eagle-mini-cpmfor-causal-lm` | `EagleMiniCPMForCausalLM` | `registry.py:590`; `vllm/model_executor/models/minicpm_eagle.py::EagleMiniCPMForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; EAGLE/EAGLE3 | ☐ required | `INVENTORIED` | none | unassigned |
-| `MODEL-SPEC-qwen3-dflash-dflash-qwen3-for-causal-lm` | `DFlashDraftModel` | `registry.py:591`; `vllm/model_executor/models/qwen3_dflash.py::DFlashQwen3ForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; sliding-window attention; DFlash | ✅ [DFlash spec](specs/dflash-spec-decode.md) §0 D14 | `DONE` | D0-redo (`CLAIM-DFLASH-D0D1`, 2026-07-26) UNBLOCKED on the advanced pin `555967922`/vLLM 0.26.0.dev0: the mixed-SWA/full z-lab 27B draft CONSTRUCTS + drafter ALIVE under `VLLM_USE_V2_MODEL_RUNNER=1` (acceptance > 1); D0+D1 (`DF-AUX-TAPS` target-side multi-tap) landed. **D2 (`DF-DRAFT-MODEL`, `CLAIM-DFLASH-D2`, 2026-07-26) CODE LANDED + CPU-GATED:** drafter model `qwen3_dflash.{h,cpp}` (plain 5-layer Qwen3-dense reusing the landed dense block ops, routing attention through the NEW `vt::DFlashBlockAttention` non-causal in-block primitive), the fc aux-combine, mask-embed, per-layer SWA/full resolution (`ResolveQwen3DFlashAttnModes`), and the z-lab loader. CPU gate GREEN: `test_ops_dflash_block_attn` 5/12 (hand-checked non-causal + RED causal-vs-non-causal + block isolation + SWA window + GQA), `test_qwen3_dflash_forward` 5/95 (forward runs + RED full-layer-causal-flip + block isolation + fc-combine ref + RED reversed-tap-order + attn-mode resolution). Existing causal `test_ops_attention` 9/9·23 + `test_qwen3_forward` 5/1028 UNCHANGED (new op is SEPARATE; causal path byte-identical). **`ACTIVE` — GPU promotion GREEN on dgx 2026-07-26 (GB10 sm_121a):** CUDA `-Werror` build clean (kernel compiles as-written); CUDA==CPU parity `test_ops_dflash_block_attn` 198412/198412 + compute-sanitizer 0; draft-forward parity `test_qwen3_dflash_draft_parity` vs the REAL vLLM DFlash draft (via `collective_rpc`) — fc rel-L2 0.46%, per-layer hidden ≤1.30%, final 0.88%, 11 deterministic rows STRICT top-1 + 5 bf16-near-tie cluster-matched; 27B SACRED 235/235 + MTP 9/9 byte-identical. Loader now tolerates the draft ckpt's omitted embed_tokens/lm_head (target-shared). Capture harness `scripts/spec/d2_dflash_draft_ref.py`. **D3 `DF-DRAFT-KV-PREP` CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D3`):** `PrepareDflashInputs` (pure-integer HOST port of `_prepare_dflash_inputs_kernel`) + `PrecomputeContextKV` + `ForwardBlockLogitsWithContext` in `qwen3_dflash.{h,cpp}` reuse the landed MatmulBT/RmsNorm/RopeNeox + the UNCHANGED D2 `vt::DFlashBlockAttention` via a `[context;block]` combined sequence (NO new kernel); CPU gate `tests/vllm/v1/spec_decode/test_dflash_kvprep.cpp` 6/114 (prepare INTEGER bit-exact + RED valid_ctx_end; context-KV envelope + RED hidden_norm/k_norm/pos; context forward degenerates to D2 at empty ctx + diverges with ctx + block isolation); additive-only diff so D2/MTP/SACRED byte-identical. **D3 DONE 2026-07-26 — GPU numeric-parity GREEN on dgx (GB10 sm_121a):** CUDA `-Werror` 0 warnings; `test_qwen3_dflash_kvprep_parity` 61/61 vs the REAL vLLM draft — `prepare_dflash_inputs` INTEGER bit-exact vs vLLM's ACTUAL Triton `_prepare_dflash_inputs_kernel` (`kernel_matches_numpy`), context-KV worst K/V rel-L2 0.31%/0.26% across all 5 layers (proves per-layer `qkv_proj` slicing == vLLM's fused KV weight), block-proposal 13 STRICT + 3 near-tie = 16/16; CPU `test_dflash_kvprep` 114/114 re-passed (RED-proven); 27B SACRED 235/235 + MTP 9/9 + D2 parity 37/37 byte-identical. Capture harness `scripts/spec/d3_dflash_kvprep_ref.py`. **D4 `DF-ENGINE-INTEGRATION` propose brick + `dflash` config-select CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D4D5`):** NEW `src/vllm/v1/worker/gpu/spec_decode/dflash/speculator.{h,cpp}` `DflashProposeBlock`/`SampleDflashBlockDrafts` — the non-autoregressive whole-block propose swapping in for the MTP k=1 `MtpProposePrefill`, composing D3 `ForwardBlockLogitsWithContext` + greedy per-mask argmax (anchor not sampled, `dflash/speculator.py:300-413`); `ParseSpeculativeConfigJson`/`ResolveDflash` accept `method:"dflash"` (lookahead k+1). CPU gate `tests/vllm/v1/spec_decode/test_dflash_propose.cpp` 5/19 GREEN (RED-first: sampler reading the anchor row fails 4/5 cases; brick == forward+sampler; empty-ctx degenerates to D2 context-free argmax; config parses; dspark throws), `-Werror`-clean TU. Additive + config-gated ⇒ MTP/non-spec byte-identical BY CONSTRUCTION (no runner/model/loader/scheduler edit). **D5 `DF-ENGINE-INTEGRATION` runner-loop LANDED + e2e RUNS on dgx 2026-07-26 (`CLAIM-DFLASH-D5`):** the verify/propose loop is wired — loader loads the SEPARATE z-lab draft (`LoadDflashDraft`, host bf16 + target-SHARED bf16 embed/lm_head) via a `--speculative-config` `model` key + `ResolveSpecConfig` dflash branch + `runner.set_dflash_draft`; the verify forward captures the D1 multi-tap (`aux_tap`→`ForwardDeviceMultiTap`); `propose_drafts_dflash` ACCUMULATES the per-request combined-feature context (`CombineAuxFeatures(aux_tap)`) and honors the `num_rejected` rollback by appending only the accepted-prefix features, then runs `DflashProposeBlock` (k=16). **e2e `test_qwen27_dflash_spec_decode` (4×32 tok vs the committed vLLM-DFlash-ON golden): 2/4 STRICT token-exact (fibonacci, three-laws) + acceptance ~ vLLM on ALL 4 (accepted 19/39/29/25 vs golden 17/39/30/25 — the MANDATORY dead-drafter-trap condition MET).** The 2 divergences (France tok11, 17*23 tok12) are SINGLE bf16 near-tie flips (17*23 RE-CONVERGES = proven near-tie; France cascades from one flip) — the ratified near-tie ROOT, rooted in the D3-documented inline bf16 context-KV recompute envelope (~0.3-1.3% rel-L2), NOT a wiring bug (proven by the 2 exact prompts + near-exact acceptance). Inertness GREEN (SACRED 235/235 + MTP 9/9 byte-identical); CUDA `-Werror` clean; NO new CUDA kernel (host orchestration reusing D1/D2/D3-sanitized ops). NOT a clean strict-4/4 pass — STRICT 4/4 token-identity + speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused context-KV projections). **D6 `CLAIM-DFLASH-D6` 2026-07-27 (records-only, NO code) — c1 SPEED A/B DONE + STRICT proven bf16-IRREDUCIBLE + CG feasibility:** our DFlash-ON = **2.50x TPOT (40.4 vs 101.2 ms) / 2.48x output-tput (24.4 vs 9.86 tok/s)** over our OFF at c1 (8 prose+code prompts×256 tok greedy, 2 reps, acceptance 0.22 = 3.56/16, rep-stable <1.5%), `benchmark_binding=true`; vs vLLM-DFlash-ON graphed vLLM-DFlash-ON graphed = 28.5 tok/s / 35.1 ms TPOT / acceptance_len 4.30 (same 8 prompts, `VLLM_USE_V2_MODEL_RUNNER=1`, mm-off, gpu_util 0.30), so OURS IS ~14% BELOW vLLM-DFlash-ON on output throughput (24.4 vs 28.5 tok/s) - both ~on-par at spec-OFF (9.86 vs 9.83 tok/s), but vLLM extracts a larger DFlash speedup (2.90x vs our 2.47x) because its draft step is fully device-resident + CUDA-graphed (ours host-orchestrates 13 downloads/step) + slightly higher acceptance (~4.3 vs ~3.6 draft tokens/step). The DONE speed bar (ours >= vLLM) is NOT met; closing it = the device-resident draft rewrite + FULL CG (D6 part 2). STRICT-4/4 is bf16-IRREDUCIBLE — the draft KV cache is bf16 not fp8 (`torch_utils.py:398` `auto`→model dtype; the D0 "fp8-KV" was the backend name), the D3 golden already compares pre-storage bf16 (residual = sub-ULP kernel noise), and a fused multi-layer KV GEMM is per-element invariant to our per-layer GEMMs (bit-exact needs vLLM's exact kernels) ⇒ the ratified near-tie gate is the FINAL correctness form (no fused-KV code landed). FULL CG + persistent-paged-KV BLOCKED on a device-resident draft-path rewrite (the D5 path does 13 device→host downloads/step) — the remaining throughput-parity increment. Inertness by construction (D5 binary; SACRED 235/235 + MTP 9/9 stand). Row STAYS `ACTIVE` (correctness-final + c1-speed-measured). Evidence tool `scripts/spec/vllm_dflash_timing.py`. **D7 `CLAIM-DFLASH-D7` 2026-07-27 (source-owning) — within-step draft forward DEVICE-RESIDENT:** `PrecomputeContextKVDevice` keeps per-layer K/V on device + `ForwardBlockLogitsWithContext` builds `[context;block]` with `vt::IndexCopy`/`IndexSelect` (removes ~30 D→H `Download`s/step), bit-identical (identity `bf16↔f32` round-trips replaced, no float op changed) — e2e `test_qwen27_dflash_spec_decode` 27/27 SAME tokens (2/4 STRICT + 2/4 near-tie at identical divergence points, acceptance 19/39/29/25), SACRED 235/235 + MTP 9/9, CUDA `-Werror` clean, compute-sanitizer 0 (198412; no new kernel). The direct old-vs-new A/B = **+2.0% output-tput (IN-NOISE)** ⇒ **D6's "downloads = the ~14% gap" REFUTED by measurement**; ours 19.68 tok/s STILL **~33% BELOW** vLLM-DFlash-ON 29.2 (reconstructed 8-prompt set; OFF parity our 9.97 ≥ vLLM 9.66); residual re-attributed to **acceptance** (ours 2.49 vs vLLM ~3.13 accepted draft-tok/step, bf16-irreducible) + per-step **context-KV RECOMPUTE** (O(context²); needs the cross-step persistent paged draft-KV store) + eager-vs-graphed. SPEED BAR NOT met; STAYS `ACTIVE` **D9 `CLAIM-DFLASH-D9` 2026-07-27 (source-owning) — PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%):** `qwen3_dflash.{h,cpp}` `AppendContextKVHost` (project ONLY newly-accepted rows → per-layer bf16 K/V) + `ForwardBlockLogitsWithPrecomputedKV` (upload the persistent store, NO re-projection) share the core `ForwardWithCtxKVDev` with the old recompute; `runner.cpp::propose_drafts_dflash` swaps the O(context²) per-step recompute (`dflash_ctx_feats_`) for an append-only per-request `dflash_kv_store_` (rollback=don't-append). NO new CUDA kernel; config-gated. Bit-identity: CPU `test_dflash_propose` 2 new D9 cases = exact float equality vs full recompute; GPU e2e `test_qwen27_dflash_spec_decode` **27/27 SAME tokens** (acceptance 19/39/29/25, same divergences); SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean. **A/B (c1, 8 prose+code×256 tok, `benchmark_binding=true`):** ours-ON **25.75 tok/s** (was 20.99, +22.7%) / 38.40 ms / acc 3.68/step vs vLLM-ON graphed **28.09** / 35.60 / acc 3.31 = **0.917×** (was 0.69×). **Part 1:** same-trajectory per-step acceptance == vLLM EXACTLY (ratio 1.00 on the 2 token-identical prompts) AND ours realized acceptance (3.68) > vLLM (3.31) ⇒ D8's "bf16 acceptance ceiling" REFUTED (confound). Residual ~8% = eager-vs-graphed ONLY; the FULL uniform-(1+k) CG (device paged-KV store + paged attn, new-CUDA) is the SOLE un-landed increment. **D12/D13 `CLAIM-DFLASH-D12/D13` 2026-07-27 — fixed-capacity paged draft-KV store (`vt::DFlashPagedBlockAttention`) + the draft-step CUDA graph landed (capture-correctness PROVEN replayed==eager bit-identical); c1 0.917×→0.978×.** **D14 `CLAIM-DFLASH-D14` 2026-07-27 — SPEED GATE MET → `DONE`:** nsys attributed the D13 ~2% residual to the from-scratch `DFlashPagedBlockAttentionKernel` draft attention (1.8% of GPU time, ~460 us/call; the draft bf16 GEMMs run `cutlass_80_wmma` in BOTH engines, NOT the gap); ported it to a WARP-scoped `__shfl_xor` online-softmax variant (mirrors the shipped `AttentionWarpKernel`; `src/vt/cuda/cuda_ops.cu` `DFlashPagedBlockAttentionWarpKernel`, default ON, `VT_DFLASH_ATTN_BLOCK=1` = the bit-identical D12/D13 block kernel) → draft attn 242.9→77.9 ms (3.1×), our-ON c1 28.60→29.32 tok/s. FINAL same-session 3-rep A/B: our-ON 29.42/29.27/29.32 vs vLLM-ON 29.240/29.247/29.233 — our WORST > vLLM's BEST, 1.003× ⇒ **≥vLLM MET.** Correctness UNCHANGED (e2e 27/27 graph==eager, acceptance 19/39/29/25 identical, 1629 draft accepted identical warp-vs-block; CUDA==CPU 795648/795648 + compute-sanitizer 0), inertness SACRED 235/235 + MTP 9/9. Correctness-complete (ratified near-tie) AND at/above vLLM throughput. Anchors: [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu#L1433) + [test_ops_dflash_paged_block_attn](../tests/vt/test_ops_dflash_paged_block_attn.cpp#L79) + [ledger](parity-ledger.md#L738); closing commit `164453a2` (claims CLAIM-DFLASH-D0D1..D14 recorded in coordination.md + ledger). | `164453a2` |
+| `MODEL-SPEC-qwen3-dflash-dflash-qwen3-for-causal-lm` | `DFlashDraftModel` | `registry.py:591`; `vllm/model_executor/models/qwen3_dflash.py::DFlashQwen3ForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; sliding-window attention; DFlash | ✅ [DFlash spec](specs/dflash-spec-decode.md) §0 D14 | `DONE` | D0-redo (`CLAIM-DFLASH-D0D1`, 2026-07-26) UNBLOCKED on the advanced pin `555967922`/vLLM 0.26.0.dev0: the mixed-SWA/full z-lab 27B draft CONSTRUCTS + drafter ALIVE under `VLLM_USE_V2_MODEL_RUNNER=1` (acceptance > 1); D0+D1 (`DF-AUX-TAPS` target-side multi-tap) landed. **D2 (`DF-DRAFT-MODEL`, `CLAIM-DFLASH-D2`, 2026-07-26) CODE LANDED + CPU-GATED:** drafter model `qwen3_dflash.{h,cpp}` (plain 5-layer Qwen3-dense reusing the landed dense block ops, routing attention through the NEW `vt::DFlashBlockAttention` non-causal in-block primitive), the fc aux-combine, mask-embed, per-layer SWA/full resolution (`ResolveQwen3DFlashAttnModes`), and the z-lab loader. CPU gate GREEN: `test_ops_dflash_block_attn` 5/12 (hand-checked non-causal + RED causal-vs-non-causal + block isolation + SWA window + GQA), `test_qwen3_dflash_forward` 5/95 (forward runs + RED full-layer-causal-flip + block isolation + fc-combine ref + RED reversed-tap-order + attn-mode resolution). Existing causal `test_ops_attention` 9/9·23 + `test_qwen3_forward` 5/1028 UNCHANGED (new op is SEPARATE; causal path byte-identical). **`ACTIVE` — GPU promotion GREEN on dgx 2026-07-26 (GB10 sm_121a):** CUDA `-Werror` build clean (kernel compiles as-written); CUDA==CPU parity `test_ops_dflash_block_attn` 198412/198412 + compute-sanitizer 0; draft-forward parity `test_qwen3_dflash_draft_parity` vs the REAL vLLM DFlash draft (via `collective_rpc`) — fc rel-L2 0.46%, per-layer hidden ≤1.30%, final 0.88%, 11 deterministic rows STRICT top-1 + 5 bf16-near-tie cluster-matched; 27B SACRED 235/235 + MTP 9/9 byte-identical. Loader now tolerates the draft ckpt's omitted embed_tokens/lm_head (target-shared). Capture harness `scripts/spec/d2_dflash_draft_ref.py`. **D3 `DF-DRAFT-KV-PREP` CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D3`):** `PrepareDflashInputs` (pure-integer HOST port of `_prepare_dflash_inputs_kernel`) + `PrecomputeContextKV` + `ForwardBlockLogitsWithContext` in `qwen3_dflash.{h,cpp}` reuse the landed MatmulBT/RmsNorm/RopeNeox + the UNCHANGED D2 `vt::DFlashBlockAttention` via a `[context;block]` combined sequence (NO new kernel); CPU gate `tests/vllm/v1/spec_decode/test_dflash_kvprep.cpp` 6/114 (prepare INTEGER bit-exact + RED valid_ctx_end; context-KV envelope + RED hidden_norm/k_norm/pos; context forward degenerates to D2 at empty ctx + diverges with ctx + block isolation); additive-only diff so D2/MTP/SACRED byte-identical. **D3 DONE 2026-07-26 — GPU numeric-parity GREEN on dgx (GB10 sm_121a):** CUDA `-Werror` 0 warnings; `test_qwen3_dflash_kvprep_parity` 61/61 vs the REAL vLLM draft — `prepare_dflash_inputs` INTEGER bit-exact vs vLLM's ACTUAL Triton `_prepare_dflash_inputs_kernel` (`kernel_matches_numpy`), context-KV worst K/V rel-L2 0.31%/0.26% across all 5 layers (proves per-layer `qkv_proj` slicing == vLLM's fused KV weight), block-proposal 13 STRICT + 3 near-tie = 16/16; CPU `test_dflash_kvprep` 114/114 re-passed (RED-proven); 27B SACRED 235/235 + MTP 9/9 + D2 parity 37/37 byte-identical. Capture harness `scripts/spec/d3_dflash_kvprep_ref.py`. **D4 `DF-ENGINE-INTEGRATION` propose brick + `dflash` config-select CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D4D5`):** NEW `src/vllm/v1/worker/gpu/spec_decode/dflash/speculator.{h,cpp}` `DflashProposeBlock`/`SampleDflashBlockDrafts` — the non-autoregressive whole-block propose swapping in for the MTP k=1 `MtpProposePrefill`, composing D3 `ForwardBlockLogitsWithContext` + greedy per-mask argmax (anchor not sampled, `dflash/speculator.py:300-413`); `ParseSpeculativeConfigJson`/`ResolveDflash` accept `method:"dflash"` (lookahead k+1). CPU gate `tests/vllm/v1/spec_decode/test_dflash_propose.cpp` 5/19 GREEN (RED-first: sampler reading the anchor row fails 4/5 cases; brick == forward+sampler; empty-ctx degenerates to D2 context-free argmax; config parses; dspark throws), `-Werror`-clean TU. Additive + config-gated ⇒ MTP/non-spec byte-identical BY CONSTRUCTION (no runner/model/loader/scheduler edit). **D5 `DF-ENGINE-INTEGRATION` runner-loop LANDED + e2e RUNS on dgx 2026-07-26 (`CLAIM-DFLASH-D5`):** the verify/propose loop is wired — loader loads the SEPARATE z-lab draft (`LoadDflashDraft`, host bf16 + target-SHARED bf16 embed/lm_head) via a `--speculative-config` `model` key + `ResolveSpecConfig` dflash branch + `runner.set_dflash_draft`; the verify forward captures the D1 multi-tap (`aux_tap`→`ForwardDeviceMultiTap`); `propose_drafts_dflash` ACCUMULATES the per-request combined-feature context (`CombineAuxFeatures(aux_tap)`) and honors the `num_rejected` rollback by appending only the accepted-prefix features, then runs `DflashProposeBlock` (k=16). **e2e `test_qwen27_dflash_spec_decode` (4×32 tok vs the committed vLLM-DFlash-ON golden): 2/4 STRICT token-exact (fibonacci, three-laws) + acceptance ~ vLLM on ALL 4 (accepted 19/39/29/25 vs golden 17/39/30/25 — the MANDATORY dead-drafter-trap condition MET).** The 2 divergences (France tok11, 17*23 tok12) are SINGLE bf16 near-tie flips (17*23 RE-CONVERGES = proven near-tie; France cascades from one flip) — the ratified near-tie ROOT, rooted in the D3-documented inline bf16 context-KV recompute envelope (~0.3-1.3% rel-L2), NOT a wiring bug (proven by the 2 exact prompts + near-exact acceptance). Inertness GREEN (SACRED 235/235 + MTP 9/9 byte-identical); CUDA `-Werror` clean; NO new CUDA kernel (host orchestration reusing D1/D2/D3-sanitized ops). NOT a clean strict-4/4 pass — STRICT 4/4 token-identity + speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused context-KV projections). **D6 `CLAIM-DFLASH-D6` 2026-07-27 (records-only, NO code) — c1 SPEED A/B DONE + STRICT proven bf16-IRREDUCIBLE + CG feasibility:** our DFlash-ON = **2.50x TPOT (40.4 vs 101.2 ms) / 2.48x output-tput (24.4 vs 9.86 tok/s)** over our OFF at c1 (8 prose+code prompts×256 tok greedy, 2 reps, acceptance 0.22 = 3.56/16, rep-stable <1.5%), `benchmark_binding=true`; vs vLLM-DFlash-ON graphed vLLM-DFlash-ON graphed = 28.5 tok/s / 35.1 ms TPOT / acceptance_len 4.30 (same 8 prompts, `VLLM_USE_V2_MODEL_RUNNER=1`, mm-off, gpu_util 0.30), so OURS IS ~14% BELOW vLLM-DFlash-ON on output throughput (24.4 vs 28.5 tok/s) - both ~on-par at spec-OFF (9.86 vs 9.83 tok/s), but vLLM extracts a larger DFlash speedup (2.90x vs our 2.47x) because its draft step is fully device-resident + CUDA-graphed (ours host-orchestrates 13 downloads/step) + slightly higher acceptance (~4.3 vs ~3.6 draft tokens/step). The DONE speed bar (ours >= vLLM) is NOT met; closing it = the device-resident draft rewrite + FULL CG (D6 part 2). STRICT-4/4 is bf16-IRREDUCIBLE — the draft KV cache is bf16 not fp8 (`torch_utils.py:398` `auto`→model dtype; the D0 "fp8-KV" was the backend name), the D3 golden already compares pre-storage bf16 (residual = sub-ULP kernel noise), and a fused multi-layer KV GEMM is per-element invariant to our per-layer GEMMs (bit-exact needs vLLM's exact kernels) ⇒ the ratified near-tie gate is the FINAL correctness form (no fused-KV code landed). FULL CG + persistent-paged-KV BLOCKED on a device-resident draft-path rewrite (the D5 path does 13 device→host downloads/step) — the remaining throughput-parity increment. Inertness by construction (D5 binary; SACRED 235/235 + MTP 9/9 stand). Row STAYS `ACTIVE` (correctness-final + c1-speed-measured). Evidence tool `scripts/spec/vllm_dflash_timing.py`. **D7 `CLAIM-DFLASH-D7` 2026-07-27 (source-owning) — within-step draft forward DEVICE-RESIDENT:** `PrecomputeContextKVDevice` keeps per-layer K/V on device + `ForwardBlockLogitsWithContext` builds `[context;block]` with `vt::IndexCopy`/`IndexSelect` (removes ~30 D→H `Download`s/step), bit-identical (identity `bf16↔f32` round-trips replaced, no float op changed) — e2e `test_qwen27_dflash_spec_decode` 27/27 SAME tokens (2/4 STRICT + 2/4 near-tie at identical divergence points, acceptance 19/39/29/25), SACRED 235/235 + MTP 9/9, CUDA `-Werror` clean, compute-sanitizer 0 (198412; no new kernel). The direct old-vs-new A/B = **+2.0% output-tput (IN-NOISE)** ⇒ **D6's "downloads = the ~14% gap" REFUTED by measurement**; ours 19.68 tok/s STILL **~33% BELOW** vLLM-DFlash-ON 29.2 (reconstructed 8-prompt set; OFF parity our 9.97 ≥ vLLM 9.66); residual re-attributed to **acceptance** (ours 2.49 vs vLLM ~3.13 accepted draft-tok/step, bf16-irreducible) + per-step **context-KV RECOMPUTE** (O(context²); needs the cross-step persistent paged draft-KV store) + eager-vs-graphed. SPEED BAR NOT met; STAYS `ACTIVE` **D9 `CLAIM-DFLASH-D9` 2026-07-27 (source-owning) — PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%):** `qwen3_dflash.{h,cpp}` `AppendContextKVHost` (project ONLY newly-accepted rows → per-layer bf16 K/V) + `ForwardBlockLogitsWithPrecomputedKV` (upload the persistent store, NO re-projection) share the core `ForwardWithCtxKVDev` with the old recompute; `runner.cpp::propose_drafts_dflash` swaps the O(context²) per-step recompute (`dflash_ctx_feats_`) for an append-only per-request `dflash_kv_store_` (rollback=don't-append). NO new CUDA kernel; config-gated. Bit-identity: CPU `test_dflash_propose` 2 new D9 cases = exact float equality vs full recompute; GPU e2e `test_qwen27_dflash_spec_decode` **27/27 SAME tokens** (acceptance 19/39/29/25, same divergences); SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean. **A/B (c1, 8 prose+code×256 tok, `benchmark_binding=true`):** ours-ON **25.75 tok/s** (was 20.99, +22.7%) / 38.40 ms / acc 3.68/step vs vLLM-ON graphed **28.09** / 35.60 / acc 3.31 = **0.917×** (was 0.69×). **Part 1:** same-trajectory per-step acceptance == vLLM EXACTLY (ratio 1.00 on the 2 token-identical prompts) AND ours realized acceptance (3.68) > vLLM (3.31) ⇒ D8's "bf16 acceptance ceiling" REFUTED (confound). Residual ~8% = eager-vs-graphed ONLY; the FULL uniform-(1+k) CG (device paged-KV store + paged attn, new-CUDA) is the SOLE un-landed increment. **D12/D13 `CLAIM-DFLASH-D12/D13` 2026-07-27 — fixed-capacity paged draft-KV store (`vt::DFlashPagedBlockAttention`) + the draft-step CUDA graph landed (capture-correctness PROVEN replayed==eager bit-identical); c1 0.917×→0.978×.** **D14 `CLAIM-DFLASH-D14` 2026-07-27 — SPEED GATE MET → `DONE`:** nsys attributed the D13 ~2% residual to the from-scratch `DFlashPagedBlockAttentionKernel` draft attention (1.8% of GPU time, ~460 us/call; the draft bf16 GEMMs run `cutlass_80_wmma` in BOTH engines, NOT the gap); ported it to a WARP-scoped `__shfl_xor` online-softmax variant (mirrors the shipped `AttentionWarpKernel`; `src/vt/cuda/cuda_ops.cu` `DFlashPagedBlockAttentionWarpKernel`, default ON, `VT_DFLASH_ATTN_BLOCK=1` = the bit-identical D12/D13 block kernel) → draft attn 242.9→77.9 ms (3.1×), our-ON c1 28.60→29.32 tok/s. FINAL same-session 3-rep A/B: our-ON 29.42/29.27/29.32 vs vLLM-ON 29.240/29.247/29.233 — our WORST > vLLM's BEST, 1.003× ⇒ **≥vLLM MET.** Correctness UNCHANGED (e2e 27/27 graph==eager, acceptance 19/39/29/25 identical, 1629 draft accepted identical warp-vs-block; CUDA==CPU 795648/795648 + compute-sanitizer 0), inertness SACRED 235/235 + MTP 9/9. Correctness-complete (ratified near-tie) AND at/above vLLM throughput. Anchors: [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu#L1433) + [test_ops_dflash_paged_block_attn](../tests/vt/test_ops_dflash_paged_block_attn.cpp#L79) + [ledger](parity-ledger.md#L738); closing commit `164453a2` (claims CLAIM-DFLASH-D0D1..D14 recorded in coordination.md + ledger). | `489a7544` |
 | `MODEL-SPEC-deepseek-v4-dspark-deepseek-v4-for-causal-lm` | `DSparkDraftModel` | `registry.py:592`; `vllm/models/deepseek_v4/__init__.py::DSparkDeepseekV4ForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; FusedMoE/grouped GEMM; MLA/latent KV; sliding-window attention; DSpark | ☐ required | `INVENTORIED` | none | unassigned |
 | `MODEL-SPEC-qwen3-dspark-qwen3-dspark-for-causal-lm` | `Qwen3DSparkModel` | `registry.py:593`; `vllm/model_executor/models/qwen3_dspark.py::Qwen3DSparkForCausalLM` | speculative draft / target-dependent | draft runner; acceptance/sampling; DSpark | ☐ required | `INVENTORIED` | none | unassigned |
 | `MODEL-SPEC-laguna-dflash-dflash-laguna-for-causal-lm` | `DFlashLagunaForCausalLM` (v0.25.0 target-pending) | v0.25.0 target `registry.py:598`; `vllm/model_executor/models/laguna_dflash.py::DFlashLagunaForCausalLM` @ `702f481` | speculative draft / Laguna targets | draft runner; acceptance/sampling; full/sliding attention; DFlash | ☐ required | `INVENTORIED` | none | unassigned |
diff --git a/.agents/state.md b/.agents/state.md
index 0fc2d622..7dfc8423 100644
--- a/.agents/state.md
+++ b/.agents/state.md
@@ -34430,3 +34430,5 @@ PARTIAL on two further unrelated test-only GCC 12 `-Wrestrict` diagnostics in
 `test_deepseek_v2_paged_engine.cpp` (actively MLA-owned) and
 `test_glm4_moe_lite_paged_engine.cpp`; those files were not touched, so full
 CTest cannot run from this incomplete build. No GPU, model, or benchmark ran.
+
+**Canonical DONE-owner reachability repaired (2026-07-31, `CLAIM-RECORD-DONE-OWNER-REPAIR`).** Records-only correction: the six checker-reported `DONE` owner cells now name the reachable commits that introduced their binding closures (`KV-PREFIX-CACHE` `a41af480`, `SAMPLE-LOGPROBS` `ecda3ce1`, and the four DFlash rows `489a7544`). Lifecycle states, evidence, support claims, source, tests, and benchmark results are unchanged.
diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md
index 2591d162..58fe7506 100644
--- a/docs/BENCHMARKS.md
+++ b/docs/BENCHMARKS.md
@@ -10967,3 +10967,10 @@ diagnostics outside this diff. The row remains
 `ANCHOR-BACKFILL`; W1-W5 in
 [the spec](../.agents/specs/c-api-library.md) name the remaining compatibility,
 no-throw, lifetime, and standalone-consumer gates.
+
+## Canonical record-owner repair (2026-07-31) - NOT APPLICABLE
+
+Six existing `DONE` matrix rows now point to reachable commits that introduced
+their binding closure states. This records-only metadata repair changes no
+lifecycle state, implementation, workload, or accepted benchmark result, so no
+new performance number is applicable or claimed.
diff --git a/docs/STATUS.md b/docs/STATUS.md
index 07739537..dccfee08 100644
--- a/docs/STATUS.md
+++ b/docs/STATUS.md
@@ -2059,3 +2059,5 @@ _(Laguna W11 go/no-go: decode is GPU-COMPUTE-bound (GPU-busy 2.56s ≈ sync 2.59
 **qwen3_5 A3 W3b LANDED (2026-07-31) — A3 COMPLETE.** Per-expert keep-quant MoE loop → 3 grouped vt::MatmulBTQuantGrouped launches (VT_QWEN35_GROUPED_MOE default-ON). BYTE-EXACT: DGX APEX-Compact grouped(=1) vs per-expert(=0) continuations byte-identical + strict-passes golden. Shared descriptor with Laguna W9. See spec §W3b.
 
 **Fold-plan A3 tier CLOSED (2026-07-31).** Both keep-quant models (Laguna W9 + qwen3_5 GGUF W2/W3a/W3b) route their routed-expert MoE through the shared vt::MatmulBTQuantGrouped descriptor, byte-exact gated. See .agents/specs/arch-fusion-fold-plan-2026-07-30.md A3.
+
+**Canonical record-owner repair (2026-07-31).** No lifecycle or support status changed. Six existing `DONE` rows now point to their reachable binding closure commits; this is records-only metadata maintenance.