diff --git a/.agents/issue-index.md b/.agents/issue-index.md
index 1b1019aad..5b1b1f17d 100644
--- a/.agents/issue-index.md
+++ b/.agents/issue-index.md
@@ -595,6 +595,7 @@ rather than merged. `scripts/check-agent-record.py` gates both.
 | [#1659](https://github.com/mudler/vllm.cpp/issues/1659) | `SPEC-DFLASH2` | **The speed harness asserts `--attention-backend` and never passes it to `LLM()`, so the declared denominator is unreachable by any path.** Measured on `dgx:gpu0` 2026-08-22 at `bed3feae6`. `capture()` built `LLM(...)` with no backend kwarg while `attention_backend_reasons` requires `resolved == declared`, and the arm log under a declared `TRITON_ATTN` reads `Using FLASH_ATTN attention backend out of potential backends: ['FLASH_ATTN','FLASHINFER','TRITON_ATTN','FLEX_ATTENTION']`. So repairing [#1658](https://github.com/mudler/vllm.cpp/issues/1658) ALONE would make the arm resolve `FLASH_ATTN`, compare it against the declared `TRITON_ATTN` and refuse on the mismatch: a working gate needs this, the probe entry and [#1657](https://github.com/mudler/vllm.cpp/issues/1657), not any one of them. This reproduces W6's recorded failure at a later head -- [#1456](https://github.com/mudler/vllm.cpp/issues/1456) measured vLLM's vendored flash-attention unable to target sm_12x, which is why `TRITON_ATTN` is the declared oracle backend on this box at all. FIXED IN FLOW: the declared backend is passed over the spellings `ATTENTION_BACKEND_KWARGS` names -- `attention_config` carrying a `backend` key first, because the measured read-back walk is `vllm_config.attention_config.backend`, then a bare `attention_backend`. The spelling is UNVERIFIED at the beyond-pin head, on the same footing `BACKEND_PROBES` was on: a wheel REJECTS a kwarg it does not declare with a `TypeError` raised while `EngineArgs` is built and therefore before anything loads, so trying both costs no lease time; a wheel that takes NEITHER is a loud refusal naming both; a spelling that is accepted and IGNORED is caught by the read-back, which is this refusal; and `--attention-backend-kwarg` pins the answer once known, with no code change. Owned by [`dflash2-spec-decode.md`](specs/dflash2-spec-decode.md) `## Owed` O28 | bug |
 | [#1660](https://github.com/mudler/vllm.cpp/issues/1660) | `SPEC-DFLASH2` | **A DFlash2 lease needs `cuda-libraries-dev-13-0` and `python3-dev`, and `RC_LEASE_ID` does not exist on this fleet.** Measured on `dgx:gpu0` 2026-08-22 over leases `52ac5673` and `a03f34e4`. DFlash2's `compute_candidates` -> `_topk` -> `flashinfer.topk` JIT-compiles `topk.cu`, which includes `<curand.h>`; the `cuda-toolkit-13-0` metapackage does NOT install that header, and leg B died on it INSIDE `profile_run` **after a 12-minute model load**, presenting as a model failure rather than a missing header. Leg C installed `cuda-libraries-dev-13-0` and got past it. `python3-dev` is likewise required or Triton driver JIT compilation fails, and that failure surfaces as `Model architectures ['Qwen3_5ForConditionalGeneration'] failed to be inspected` -- which names the model, not the toolchain. Separately, the leased worker carries `RC_DEVICE`, `RC_JOB_ID` and `RC_TOKEN` and no `RC_LEASE_ID`, so the gate's `--lease-id` default read empty and the `## Owed` O26 recipe run VERBATIM refused with `no lease id. Claim the device with rc run/rc hold first; never ssh to a fleet box`. The refusal is correct on a missing lease; the defect is that the committed procedure could not satisfy its own gate. FIXED IN FLOW: the default takes `RC_JOB_ID` first and falls back to `RC_LEASE_ID` -- the measured variable first, the documented-but-absent one for a controller that does export it -- both packages are recorded in [`environment.md`](environment.md) beside the existing note that no CUDA toolkit is preinstalled, and the O26 recipe now names them. Owned by [`dflash2-spec-decode.md`](specs/dflash2-spec-decode.md) `## Owed` O28 | bug |
 | [#1543](https://github.com/mudler/vllm.cpp/issues/1543) | `ENG-RECORD-CONFLICT-SURFACES` | `check-public-doc-tables.py` gates the projection documents for size and shape only, never for key uniqueness, so `docs/STATUS.md`, `docs/BENCHMARKS.md` and `docs/FEATURES.md` can each carry two contradictory rows for one key with nothing indicating which is current. MEASURED on [PR #1081](https://github.com/mudler/vllm.cpp/pull/1081): `docs/BENCHMARKS.md` carried the key `**NemotronH paged forward** (MODEL-NEMOTRON-H-ABI-A2P, #810)` TWICE and the two rows contradict -- one says the GB10 read was 4/24 with the sm_121a re-run pending a lease, the other says the same gate PASSES 96/96 -- and the checker returned rc=0 on that file. Counted directly, the key appears once on `origin/main` and twice on the PR head. The blindness is PROVED BY MUTATION rather than inferred from reading: a reviewer reinserted a duplicate key with contradictory content and `benchmarks_errors()` still returned `[]`, and the checker holds no uniqueness or duplicate-key logic at all (positive control: 14 `errors.append` calls in the same file, so the grep pattern finds what is there). No checker asserts uniqueness ON A PROJECTION DOCUMENT, and that is the whole gap -- duplicate-key checks themselves are neither scarce nor novel here, and the closest is prior art the fixing row should mirror rather than reinvent. `scripts/check-agent-record.py::check_matrices` already refuses a `duplicate ID` over the matrix inventories, keyed on the row ID and reported against the file and line that repeats it, and `scripts/check-agent-record.py::check_issue_index` refuses `issue #N listed twice` -- the refusal this very row had to be checked against before it was appended. That same file READS `docs/BENCHMARKS.md`, through `scripts/check-agent-record.py::markdown_files`, but never for key uniqueness. The remaining duplicate-key checks are structural rather than record-keyed and transfer nothing: a duplicate YAML key in `check-release-workflow.py`, a duplicate JSON key in `ready-for-helper.py`, a duplicated oracle-pin key in `check-oracle-pins.py`. The cause is that two record surfaces carry OPPOSITE disciplines and nothing tells an author which one is open: `.agents/issue-index.md` is an append-only log where appending is correct, and `docs/BENCHMARKS.md` is a keyed projection where appending is the defect. Two review passes and one operator verification missed it for one reason worth recording -- each checked that both sides' hunks survived the merge, which proves the merge was faithful and says nothing about whether the RESULT is key-unique. SCOPE for the fixing row, so it does not over-promise: assert uniqueness on the first table cell and scope it PER TABLE, because two tables may legitimately share a key and because `audit-live-rows --check` already reports three IDs live in two matrices (BACKEND-CPU, BACKEND-CUDA-SM121, BACKEND-VULKAN), so a naive global check reds `main` on day one -- the failure this repository already suffered twice in one week in [#1464](https://github.com/mudler/vllm.cpp/issues/1464) and [#1468](https://github.com/mudler/vllm.cpp/issues/1468). The red-before is available and is not hypothetical: PR #1081's own pre-repair head contains a real duplicate. NOT fixed in flow, and this index row does not fix it -- a checker-semantics change owes its own row, spec, red-before test or mutation, and green-after evidence, and #1081's separate repair collapses its two rows into one WITHOUT closing this, because the next append reproduces it. Owned here because this row is the standing owner of public-document projection-gate semantics ([#460](https://github.com/mudler/vllm.cpp/issues/460), [#495](https://github.com/mudler/vllm.cpp/issues/495), [#498](https://github.com/mudler/vllm.cpp/issues/498), [#507](https://github.com/mudler/vllm.cpp/issues/507), [#1055](https://github.com/mudler/vllm.cpp/issues/1055)) -- named at that width deliberately, because the five do NOT all sit in one script: #460's `check_links` and #498's retired `MAX_README_CHARS` live in `check-agent-record.py` and `check-readme-structure.py`, so calling the group `check-public-doc-tables.py` semantics would be false of two of them. The row is `READY` rather than `DONE`, and its spec already names `tests/scripts/test_check_public_doc_tables.py` as a red-before suite, so the owner is a live row whose declared test surface is the checker this issue indicts; the fixing row still owes the spec and the red-before, which ownership records rather than replaces | bug |
+| [#1663](https://github.com/mudler/vllm.cpp/issues/1663) | `KERNEL-ATTN-DENSE-FLASH` | **Three stems outlived their removing rows on `scripts/attention-rung-allowlist.txt`, and while a stem sits there it excuses its WHOLE translation unit -- including a marker somebody deletes.** `47a918d8f` ([#1579](https://github.com/mudler/vllm.cpp/pull/1579), issue [#1545](https://github.com/mudler/vllm.cpp/issues/1545)) routed `muse_glimmer_vision`'s perception encoder to `vt::AttentionDenseFlash`, and `90e8c3c85` ([#1557](https://github.com/mudler/vllm.cpp/pull/1557), issue [#1549](https://github.com/mudler/vllm.cpp/issues/1549)) swapped the LTX-2.5 DiT device forward to the same op; after both, `muse_glimmer_vision` names `vt::Attention` nowhere, and `ltx2`/`ltx2_device` name it only at calls that now carry their own `VT-ATTN-NAIVE:` reason (the host CPU-only arm, where `kAttention` and `kAttentionDenseFlash` are the same registered function, and the `VLLM_LTX2_DIT_FLASH_ATTN=0` arm of a same-binary A/B). `scripts/check-attention-rung-consistency.py` reported all three `STALE (not a failure)` and exited 0, which is the header's design -- the removing row is not forced to edit this file and the deletion falls to the next preflight -- so the entries were owed rather than wrong. **The hole is MEASURED, not inferred.** On `db648fb88`, deleting the `// VT-ATTN-NAIVE:` marker at `src/vllm/model_executor/models/ltx2.cpp:959` leaves the checker at **rc=0** (`7 carry a recorded reason, 1 unmarked and excused`) with the three stems listed, and at **rc=1** naming `ltx2.cpp:966` with them removed; both arms restored byte-for-byte against a pre-taken sha256. The window was open from `90e8c3c85` until this fix. FIXED IN FLOW: the stems are deleted and `tests/scripts/test_check_attention_rung_consistency.py::ShippedTreeTests::test_allowlist_holds_only_the_in_flight_stems` moves to the empty set in the same change, as the allowlist header requires. The empty set is not a weaker assertion -- `drift_sites` now excuses nothing, so `test_shipped_tree_is_green` measures the tree on its markers alone -- and a new case asserts the three stems positively and ASYMMETRICALLY, `ltx2`/`ltx2_device` present-and-marked against `muse_glimmer_vision` absent, because one assertion covering all three is false of one of them in either direction. Distinct from [#1629](https://github.com/mudler/vllm.cpp/issues/1629), which is the population floor in the same checker and is untouched here. Spec [`attention-rung-visibility.md`](specs/attention-rung-visibility.md), parent [#1544](https://github.com/mudler/vllm.cpp/issues/1544) | bug |
 | [#1470](https://github.com/mudler/vllm.cpp/issues/1470) | `LTX25-PHASE-RESIDUE` | **`test_ltx2_video` false-redded once on `main` under load and the failing case's identity was never captured.** The prediction it records (F12 of #1441's fifth review) was that a coverage floor whose per-boundary sampling cost is fixed while the leaf shrinks is a false-RED risk and never a false pass. Measured by `LTX25-PHASE-RESIDUE`: the prediction is right about the polarity and WRONG about the term. The `denoise` coverage miss is not sampling cost, it is the sampler's post-process and Euler step, which no anchor wrapped and which scale with the latent -- 49 us per step at nine frames against 343 us at 81, in ONE run of one binary, and instrument cost does not move 7x with the latent. **STILL OPEN, and the correction above is the whole of what this row establishes.** The row's branch anchored that work as `denoise.update` and [#1556](https://github.com/mudler/vllm.cpp/issues/1556) is CLOSED rather than merged, so nothing anchors it on `main` today; the anchor is owed under [#1668](https://github.com/mudler/vllm.cpp/issues/1668). An earlier draft of this row said the issue was closed by anchoring, which was true of the branch and never of the tree. Spec [`ltx25-phase-residue.md`](specs/ltx25-phase-residue.md) | bug |
 | [#1536](https://github.com/mudler/vllm.cpp/issues/1536) | `LTX25-PHASE-RESIDUE` | **`test_ltx2_video` was persistently red rather than load-flaky, and it was the only failing test on `main`.** The issue asks for the residue to be DECOMPOSED rather than re-argued, and names `d995c52f0` (the temporal x2 upsampler) as the first hypothesis to test. The decomposition REFUTES that hypothesis and settles the cause: **92% of the un-named time is one region**, `Ltx2VideoEngine::Load` from the timeline's origin to `Open("load.dit")` -- 17.661 ms of a 19.178 ms residue -- while the upsampler's own work sits inside `phase.upsample_latent`, a named leaf that does not appear in the residue at all. The remaining four gaps are 4.95%, 1.30%, 1.09% and 0.56%, and the sixteen gaps between adjacent named phases hold 6.8 us each, which is the instrument and nothing else. **The RED is gone and the issue is NOT.** `519303d15` ([#1622](https://github.com/mudler/vllm.cpp/pull/1622), row `LTX25-DEVICE-RESIDENCY`) names that same 92% region `load.open` -- same open point, same close point, same `Scope::Close` shape as this row's `load.setup` -- so the sum floor is repaired on `main` by another row, and the coverage floor is repaired by `6b48edb2c` moving it to 0.75, which that change's own comment describes as a holding action in substance, though not in those words. What is owed is the naming that would make a tight floor honest again, under [#1668](https://github.com/mudler/vllm.cpp/issues/1668). Spec [`ltx25-phase-residue.md`](specs/ltx25-phase-residue.md) | bug |
 | [#1567](https://github.com/mudler/vllm.cpp/issues/1567) | `LTX25-PHASE-RESIDUE` | the res_2s sampler arm has no `denoise.update` anchor, so its denoise would be decomposed on one arm only. Filed by `LTX25-PHASE-RESIDUE` and deliberately NOT landed with it, because no gate in this tree renders on that arm and an anchor no gate runs is dead code. `Ltx2Res2sDenoisingLoop` runs its own post-process and step behind `Ltx2Res2sHooks`, so the anchor needs a hook rather than a statement. It lives in `ltx2_samplers.cpp`, is declared in `ltx2_samplers.h` beside the hooks struct, and is called from `ltx2_video.cpp`. **NOT `ltx2_res2s.cpp`**: #1556's spec named that file and it has never existed here, which `git log --all --diff-filter=A` confirms; #1567's forge text names no file at all, so the wrong anchor came from the spec rather than from the issue. Doubly owed while [#1668](https://github.com/mudler/vllm.cpp/issues/1668) is open, since the first-order arm has no anchor either. Listed under `## Owed` in [`ltx25-phase-residue.md`](specs/ltx25-phase-residue.md) | bug |
diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md
index 193ada943..f8545894c 100644
--- a/.agents/kernel-matrix.md
+++ b/.agents/kernel-matrix.md
@@ -150,7 +150,7 @@ host/sched. Detail: state `KERNEL-FA2-GQA-SWAP-FLIP`.
 | `KERNEL-MOE-SQRTSOFTPLUS-HASH` | **DeepSeek-V4 MoE router + clamped-SwiGLU deltas — the three genuinely-new-vs-V2/V3 MoE pieces** (DeepSeek-V4-Flash W6). V4 keeps the DeepSeek grouped-GEMM / 256-expert w13/w2 / shared-expert / NVFP4 machinery (REUSED, not re-ported) but replaces three primitives. Three ops: **(1)** the router SCORE function **`sqrt(softplus(x))`** (`softplus(x)=log(1+exp(x))`, then sqrt) — distinct from V2/V3's sigmoid/softmax `noaux_tc`; the sqrt∘softplus COMPOSITION is load-bearing (RED-first proven); **(2)** the router: score all experts, add `e_score_correction_bias` for SELECTION ONLY, pick top-k OR — for the first `num_hash_layers` HASH layers — look experts up directly in the `tid2eid` [vocab, topk] token-id→expert table (BYPASSING top-k), GATHER weights from the UNBIASED scores, renormalize, ×`routed_scaling_factor` (the bias-affects-selection-not-weights split + the hash bypass are load-bearing, both RED-first proven); **(3)** the **clamped SwiGLU** expert activation `SiluAndMulWithClamp` — `gate` clamped `max=limit` (max ONLY), `up` clamped `[-limit,+limit]` (BOTH sides), then `gate·sigmoid(alpha·gate)·(up+beta)`; the ASYMMETRIC clamp is load-bearing (RED-first proven). MegaMoE (SM100-only) is NOT the GB10 target — this mirrors the FusedMoE-fallback router GB10 runs | score `vllm/model_executor/layers/fused_moe/router/fused_topk_bias_router.py:88`; router `fused_topk_bias_router.py:75-118` (`_topk_softplus_sqrt_torch`) + hash branch `:100-106` + dispatch `:254-265`; hash-table wiring `vllm/models/deepseek_v4/nvidia/model.py:562-578,:686,:696-717`; FusedMoE fallback `nvidia/model.py:647-691`; clamped SwiGLU `vllm/model_executor/layers/activation.py:197-201` (`SiluAndMulWithClamp.forward_native`), used by DeepseekV4MLP `nvidia/model.py:126-133`; cross-checked SGLang `v0.5.15` `python/sglang/srt/layers/moe/{topk.py:1013-1014, hash_topk.py:137-180}` @ `555967922` | Portable host reference (device kernels landed W7-device — see `KERNEL-DSV4-W7-DEVICE`) [deepseek_v4_moe.cpp](../src/vllm/model_executor/models/deepseek_v4_moe.cpp) + [deepseek_v4_moe.h](../include/vllm/model_executor/models/deepseek_v4_moe.h): `SqrtSoftplus` / `SqrtSoftplusRouteTopk` / `ClampedSwiGLU` | **CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, `-Wall -Werror -Wextra` 0-warn on the new TUs):** [test_deepseek_v4_moe.cpp](../tests/vllm/models/test_deepseek_v4_moe.cpp) **12/12 cases · 716 assertions** — hand-derived literal cases (sqrt∘softplus composition `softplus(x)=4 ⇒ score=2`; bias flips selection but weight stays the UNBIASED 1.0 not 3.0; renormalize by the unbiased sum; routed_scaling_factor; hash `tid2eid` picks {3,1} where top-k would pick {2,0}; asymmetric clamp gate=-5 kept vs up clamped to -2; gate/up clamp boundaries; alpha/beta) + from-first-principles double-precision references (router f32==f64 rel-L2 < 1e-5 + exact ids; SqrtSoftplus f64 + monotonicity; ClampedSwiGLU rel-L2 < 1e-6). **RED-first PROVEN all three levers:** drop the sqrt → 8 cases/493 assertions fail; gather weights from the BIASED scores → 2 cases/181 fail; symmetric-clamp the gate → 2 cases/6 fail; revert restores 12/12·716. Honest gate form: host-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). Full-model gate multi-Spark-blocked (156.7 GiB); the device kernels reuse the existing grouped-GEMM + `DeepseekV4Model::Forward` assembly (W7) + the strict/near-tie engine gate (W8) are named residuals | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W6 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W6` |
 | `KERNEL-DSV4-W7-DEVICE` | **DeepSeek-V4-Flash W7-DEVICE — the four NEW V4 op families' CUDA kernels**, each a 1:1 DEVICE port of the landed portable HOST reference (the oracle the four SPIKE rows above pinned), registered through the vt **OpProvider seam** (`kDeepseekV4{Mhc,Dsa,Compressor,Moe}`) so `DeepseekV4Model::ForwardDevice` can dispatch them: **(MHC)** Sinkhorn + mHC pre/post + hc_head collapse; **(DSA)** indexer weight-fold + weighted-MQA ReLU logits + causal top-k + per-head attention-sink softmax + grouped output-LoRA; **(Compressor)** softmax-window pool + RMSNorm + save-time APE + **fp8_ds_mla** KV encode (UE8M0 block scale + e4m3, bf16 rope) / decode; **(MoE)** sqrtsoftplus/hash router + clamped SwiGLU. The 512-wide MLA attention + expert grouped-GEMM REUSE the existing NVFP4/FP8 kernels (`kMlaDecodeAttention`/`kMoeGroupedGemmNvfp4`, `cuda_mla_attn.cu`/`cuda_moe*.cu`) and are NOT re-ported | the SAME `file:line` the host refs cite (the `KERNEL-{MHC-SINKHORN,ATTN-DSA-SPARSE-INDEX,ATTN-DSA-COMPRESSOR,MOE-SQRTSOFTPLUS-HASH}` upstream columns), @ `555967922` | [cuda_deepseek_v4.cu](../src/vt/cuda/cuda_deepseek_v4.cu) (kernels + host-vector launchers + OpProvider registration) + [deepseek_v4_device.h](../include/vllm/model_executor/models/deepseek_v4_device.h) / [deepseek_v4_device.cpp](../src/vllm/model_executor/models/deepseek_v4_device.cpp) (seam resolvers); `DeepseekV4Model::ForwardDevice` composes them ([deepseek_v4.cpp](../src/vllm/model_executor/models/deepseek_v4.cpp)) | **DGX GB10 (sm_121a) UNIT GATE GREEN + RUNTIME-VERIFIED (2026-07-29):** [test_cuda_deepseek_v4.cpp](../tests/vllm/models/test_cuda_deepseek_v4.cpp) **11/11 cases · 153 assertions** — each device kernel vs its host-ref oracle at small shape: BIT-EXACT ids (DSA causal top-k, sqrtsoftplus/hash router selection), `-inf` mask exact (indexer out-of-window), near-tie rel-L2 < 1e-4 for the fp reductions (Sinkhorn, pool/softmax, sqrtsoftplus — device `expf`/`sqrtf`/`rsqrt` vs host), fp8_ds_mla encode→decode within the e4m3 granularity bound, bf16 rope bit-exact; PLUS the **ForwardDevice composition gate** (device forward == host forward, rel-L2 < 2e-3 over the 4-family tiny-config interleave). **compute-sanitizer memcheck 0 errors.** **RED-first PROVEN:** dropping the sqrt in the device sqrtsoftplus fails 3 cases / 6 assertions (sqrtsoftplus + router weights + ForwardDevice); revert restores 11/11·153. Build: CUDA `-Werror` clean (the #155 voxtral GCC-13 `-O2` array-bounds/stringop false positive neutralized locally). Honest 3-state: RUNTIME-VERIFIED at small shape on real GB10; the real-checkpoint paged-engine e2e stays W8 (156.7 GiB does not fit ONE GB10). **DECODE GLUE-FOLD LANDED (2026-08-03, GB10 sm_121a):** the resident-decode `norm_rope_rows` + MHC-pre kernels folded FP64→FP32 — `VT_V4_ROPE_FLOAT` (fused norm+RoPE **4.58→0.46 ms/step ~10×**, decode +6.1%) + `VT_V4_MHC_LEAN` (finish block 256→1024 + sqrsum-fold, +0.7%, floored by 86 sequential single-block launches/step), both default-ON + BYTE-EXACT (decode ids `=1`/`=0` token-identical via the resident-decode path `--gpu --kv-cache`); `test_cuda_deepseek_v4` Brick-7 + Brick-B **20/20·67073** PASS on GB10; net decode 14.02→14.96 tok/s → 90.7% of ds4 ~16.5 (`CLAIM-DSV4-ROPE-FLOAT` / `CLAIM-DSV4-MHC-LEAN`; ds4 bar corrected from the unreproduced 17.13 anchor to the fair same-session ~16.5, the later MHC-SINK4 reached ~96%) | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W7 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W7-DEVICE` |
 | `KERNEL-KDA-DELTA` | **Kimi Delta Attention (KDA) gated-linear-attention delta vs plain GDN — a genuinely new gated-linear-attention family, the shared unblocker for Kimi-Linear-48B and Kimi-K3 (W4).** `KimiGatedDeltaNetAttention` SUBCLASSES `GatedDeltaNetAttention`, so its conv-state/cache layout, `GDNAttentionMetadata`, chunked-delta recurrence and WY solve are REUSED from our landed GDN — this row owns ONLY the four KDA-specific deltas plain GDN lacks: **(1)** a per-channel **`[H,D]` low-rank decay** via an `f_a_proj→f_b_proj` bottleneck (GDN has only a per-HEAD scalar decay from `A_log`); **(2)** the decay GATE `g = -exp(A_log[h])·softplus_β(g1+dt_bias)` per channel (β=1, thr=20; `kda_gate_fwd_kernel` decode) + its chunk-local cumulative-sum prefill variant (`kda_gate_cumsum_fwd_kernel`, folds `RCP_LN2`); **(3)** the **sigmoid-gated output norm** `FusedRMSNormGated(head_dim, activation="sigmoid")` = `rmsnorm(x)·w·σ(g)` (the gated-linear-attention output norm GDN lacks); **(4)** three separate q/k/v short causal convs (`conv_size=4`, silu) + the q/k **L2-norm** preprocessing (`x/sqrt(Σx²+eps)`, SUM not mean). ADDITIVE — does NOT touch `cuda_gdn.cu`/`gdn_attn.cpp`, so the Qwen3.6-27B/35B GDN gate is structurally untouched (like DSA kept shared-MLA untouched) | decay bottleneck `vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py:142-156,:245`; decay gate `vllm/third_party/flash_linear_attention/ops/kda.py:1541-1600,:1603-1646`; chunk-cumsum `kda.py:1182-1254,:1257-1303`; gated norm `kda.py:463-487` (`:436` eps=1e-5); short conv `kimi_gdn_linear_attn.py:171-198,:324-356`; q/k L2-norm `kda.py:1511-1513` + `ops/l2norm.py:42-43,:96` @ `555967922` | Portable host reference (device kernel is a named residual) [kimi_kda.cpp](../src/vllm/model_executor/models/kimi_kda.cpp) + [kimi_kda.h](../include/vllm/model_executor/models/kimi_kda.h): `KdaLowRankDecay` / `KdaDecayGate` / `KdaDecayGateChunkCumsum` / `FusedRMSNormGated` / `KdaShortConv` / `L2NormRows` | **CPU UNIT GATE GREEN (2026-07-28, `-Wall -Werror -Wextra` 0-warn):** [test_kimi_kda.cpp](../tests/vllm/models/test_kimi_kda.cpp) **14/14 cases · 36 assertions** — hand-derived literal cases (f_b∘f_a bottleneck; `-exp(A_log)·softplus` with the >thr linearisation; per-head A_log + per-channel dt_bias; chunk-cumsum reset+`RCP_LN2` fold; sigmoid-gated norm; swish-vs-sigmoid branch; per-head-dim normalisation; causal-depthwise+silu conv; zero-init-state edge; L2-norm SUM-not-mean) + from-first-principles double-precision references on randomized shapes (decay gate, gated norm, short conv rel-L2 < 1e-6). Honest gate form: host-reference + structural review, NOT a dumped-oracle rel-L2 — the REAL e2e gate is the Kimi-Linear-48B-A3B proxy vs the pinned oracle (DGX-blocked; K3 2.8T does not fit one GB10). Named residuals: the KDA CUDA device kernel + the Kimi-Linear-48B proxy gate — anchor `tests/vllm/models/test_kimi_kda.cpp:41` | [kda-kernel-delta spike](specs/kda-kernel-delta.md) | `SPIKE` | `CLAIM-KDA-KERNEL` |
-| `KERNEL-ATTN-DENSE-FLASH` | **Flash-TILED dense non-causal attention — the SHARED-MEMORY-TILED form of `AttentionDenseFast` for long non-causal contexts** (multimodal-speed §14, the Whisper AUDIO encoder — hd-64, non-causal, 1500 frames × 32 layers). A block of `kFlashBr=16` query-warps (512 threads) SHARES each streamed `kFlashBc=64`-column K/V tile out of shared memory (classic FlashAttention K/V tiling): the CTA cooperatively loads a K/V tile into shared memory, then each warp runs its online-softmax update reading K/V from shared memory, killing `AttentionWarpKernel`'s O(t²) redundant global K/V re-reads (one full K/V sweep per (query,head)). One q-head per CTA (all warps share the GQA kv-head). BIT-IDENTICAL to `AttentionDenseFast`: the per-warp arithmetic (per-lane head_dim grouping `lane+32k`, butterfly `__shfl_xor`, sequential j-order, f32 online-softmax `m`/`l`/`acc`) is copied verbatim, only K/V bytes come from shared memory instead of global ⇒ token-identical by construction. **Head_dim-generic** (`npl=(d+31)/32`; the register blocking allows d≤256 but the K/V tile's dynamic shared memory is what BINDS — d≤192 bf16 / d≤96 f32 under CUDA's default 48 KiB cap, see the 2026-08-21 entry): since 2026-07-28 (multimodal-speed §16, `CLAIM-MM-SPEED-QWEN-IMAGE`) ALSO the default for the Qwen3-VL / Qwen3.6-27B VISION tower per-frame self-attention (hd-72, non-causal, 784 patches) — byte-identical to the warp `AttentionDenseFast` it replaced there (bench 0/1,003,520 mismatch; STRICT image/video e2e 32/32) | STRUCTURE ported 1:1 from vendored FlashAttention-2 `compute_attn_1rowblock` [flash_fwd_kernel.h:52](../src/vt/cuda/flash_attn/src/flash_fwd_kernel.h#L52) (sK/sV shared tiles :163-165 + the `for(int n_block…)` K/V-tile stream + online rescale); non-causal encoder dispatch cross-checked to vLLM `WhisperEncoderAttention` [whisper.py:255](https://github.com/vllm-project/vllm/blob/e24d1b24/vllm/model_executor/models/whisper.py#L255) | `OpId::kAttentionDenseFlash` + decl [ops.h](../include/vt/ops.h) + wrapper/validation [ops.cpp](../src/vt/ops.cpp); CUDA `AttentionDenseFlashKernel`/`AttentionDenseFlashKernelCuda` [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu); CPU maps to `AttentionKernel` (byte-identical) [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp); wired default in [whisper_audio.cpp](../src/vllm/model_executor/models/whisper_audio.cpp) (`VT_WHISPER_ENC_WARP=1`/`VT_WHISPER_ENC_EAGER=1` A/B) + [qwen3_vl_vision.cpp](../src/vllm/model_executor/models/qwen3_vl_vision.cpp) (vision tower default since §16; `VT_QWEN3VL_ATTN_WARP=1`/`VT_QWEN3VL_ATTN_EAGER=1` A/B) | **GPU GATE GREEN on dgx (2026-07-28, GB10 sm_121a, base `af1ed76b`):** CUDA `-Werror` 0-warn (cutlass-ON + FA2-ENABLED banners). `test_voxtral_e2e` **16/16** default-flash; flash/warp/eager token dumps md5-IDENTICAL (`89923566…`) ⇒ ZERO token flips; goldens md5 UNCHANGED (`voxtral_golden.json 8ab87b7e…`, `voxtral_neartie.json 937b9ad3…`, before==after). Proof-of-run nsys `AttentionDenseFlashKernel` 32 inst, ZERO `AttentionWarpKernel`/naive on encoder; RED confirmed (corrupt kernel → gate FAILS → restore → 16/16); `compute-sanitizer --tool memcheck` **0 errors**; 3 runs byte-identical. **A/B (same binary, `flock`, rep0 dropped):** attention **35.11 → 19.29 ms/layer (1.82×, NON-OVERLAPPING)**; encoder forward **~1834 → ~1375 ms (1.33×)**. **NOT at parity:** ~1.37 s vs vLLM ~43 ms TTFT (~32×, was ~44×) — the scalar warp-per-query recurrence is now serial-latency-bound over 1500 keys (L2 already served much of the redundant reads ⇒ 1.8× not 16×); gap-closer is a tensor-core MMA hd-64 non-causal FA2 instantiation (LARGE) + resident encoder weights (MEDIUM). **Vision tower (§16, 2026-07-28):** extended to the Qwen3-VL/27B tower (hd-72, 784 patches) — STRICT image/video e2e 32/32, bench flash-vs-warp 0/1,003,520 mismatch, nsys default 4B e2e `AttentionDenseFlashKernel` 24 inst/zero warp, RED 30/46→46/46, sanitizer 0; A/B warp 148.3→flash 142.3 ms = 1.04× (small — the vision attention at t=784 is serial-latency-bound not bandwidth-bound; the tower already BEATS vLLM at 0.57× eager) **2026-08-21 (`CLAIM-ATTN-RUNG-VISIBLE`, issue [#1544](https://github.com/mudler/vllm.cpp/issues/1544), spec [attention-rung-visibility.md](specs/attention-rung-visibility.md)): the advertised head_dim contract is now the LAUNCHABLE one, and the naive rung stops being a silent default.** The op stated `d <= 256` ([cuda_ops.cu](../src/vt/cuda/cuda_ops.cu) `LaunchAttentionDenseFlash`) while requesting `2*kFlashBc*d*sizeof(Tin)` bytes of DYNAMIC shared memory with no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`, so the driver's default 48 KiB cap made the real ceiling **192 bf16 / 96 f32** — Kimi (192 f32, 96 KB) and Qwen3.5 (256) would have taken a bare launch error from the `cudaGetLastError` at the bottom of the launcher, naming nothing they could do instead. The bound now lives in [ops.h](../include/vt/ops.h) as `AttentionDenseFlashSmemBytes` / `AttentionDenseFlashMaxHeadDim`, PURE host arithmetic so a box with no GPU can execute it, tied to the kernel by two `static_assert`s on `kFlashBc` and the register blocking; the launcher refuses above it naming `vt::AttentionDenseFast`, which uses NO shared memory and does serve those widths. NARROWING was chosen over `cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize)`: d=256 f32 wants 128 KiB, above the opt-in per-block cap of the consumer Blackwell parts gated here, so the opt-in would still leave the widest advertised width a lie AND cannot be verified without a device. Strictly additive for callers — the bound is INCLUSIVE, so d=192 bf16 lands exactly on 49152 and still launches. Mirrors `supports_head_size` / `get_supported_head_sizes` [backend.py:155-163](https://github.com/vllm-project/vllm/blob/555967922/vllm/v1/attention/backend.py#L155), consulted BEFORE dispatch rather than discovered by launching. Same change adds [check-attention-rung-consistency.py](../scripts/check-attention-rung-consistency.py) (preflight + CI): a model TU naming `vt::Attention` needs a `// VT-ATTN-NAIVE:` reason beside the call, so the six deliberate sites now say why and a new author gets a red instead of a silent ~500x. `kAttention` and every existing caller's numerics are UNTOUCHED by construction — the checker executes no model code and the head_dim guard only fires where the launch already failed. CPU-GATED: checker green (9 sites / 6 marked / 3 unmarked and excused by the 3 in-flight allowlisted stems), unit + mutation cases in [test_check_attention_rung_consistency.py](../tests/scripts/test_check_attention_rung_consistency.py) (no case count is recorded here -- a count of one file stored in another is a drift lock; run the suite for the live number), new head_dim contract cases in [test_ops_attention.cpp](../tests/vt/test_ops_attention.cpp). OWED [#1573](https://github.com/mudler/vllm.cpp/issues/1573): the on-device refusal case and its reachability mutation are PENDING a lease — `dgx:gpu0` was held by the developer, and the CPU cases pin the arithmetic, never that the launcher calls it. | [multimodal-speed](specs/multimodal-speed.md) §14 + §16 | `ACTIVE` | `CLAIM-MM-SPEED-AUDIO-ENC-KERNEL` + `CLAIM-MM-SPEED-QWEN-IMAGE` + `CLAIM-ATTN-RUNG-VISIBLE` |
+| `KERNEL-ATTN-DENSE-FLASH` | **Flash-TILED dense non-causal attention — the SHARED-MEMORY-TILED form of `AttentionDenseFast` for long non-causal contexts** (multimodal-speed §14, the Whisper AUDIO encoder — hd-64, non-causal, 1500 frames × 32 layers). A block of `kFlashBr=16` query-warps (512 threads) SHARES each streamed `kFlashBc=64`-column K/V tile out of shared memory (classic FlashAttention K/V tiling): the CTA cooperatively loads a K/V tile into shared memory, then each warp runs its online-softmax update reading K/V from shared memory, killing `AttentionWarpKernel`'s O(t²) redundant global K/V re-reads (one full K/V sweep per (query,head)). One q-head per CTA (all warps share the GQA kv-head). BIT-IDENTICAL to `AttentionDenseFast`: the per-warp arithmetic (per-lane head_dim grouping `lane+32k`, butterfly `__shfl_xor`, sequential j-order, f32 online-softmax `m`/`l`/`acc`) is copied verbatim, only K/V bytes come from shared memory instead of global ⇒ token-identical by construction. **Head_dim-generic** (`npl=(d+31)/32`; the register blocking allows d≤256 but the K/V tile's dynamic shared memory is what BINDS — d≤192 bf16 / d≤96 f32 under CUDA's default 48 KiB cap, see the 2026-08-21 entry): since 2026-07-28 (multimodal-speed §16, `CLAIM-MM-SPEED-QWEN-IMAGE`) ALSO the default for the Qwen3-VL / Qwen3.6-27B VISION tower per-frame self-attention (hd-72, non-causal, 784 patches) — byte-identical to the warp `AttentionDenseFast` it replaced there (bench 0/1,003,520 mismatch; STRICT image/video e2e 32/32) | STRUCTURE ported 1:1 from vendored FlashAttention-2 `compute_attn_1rowblock` [flash_fwd_kernel.h:52](../src/vt/cuda/flash_attn/src/flash_fwd_kernel.h#L52) (sK/sV shared tiles :163-165 + the `for(int n_block…)` K/V-tile stream + online rescale); non-causal encoder dispatch cross-checked to vLLM `WhisperEncoderAttention` [whisper.py:255](https://github.com/vllm-project/vllm/blob/e24d1b24/vllm/model_executor/models/whisper.py#L255) | `OpId::kAttentionDenseFlash` + decl [ops.h](../include/vt/ops.h) + wrapper/validation [ops.cpp](../src/vt/ops.cpp); CUDA `AttentionDenseFlashKernel`/`AttentionDenseFlashKernelCuda` [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu); CPU maps to `AttentionKernel` (byte-identical) [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp); wired default in [whisper_audio.cpp](../src/vllm/model_executor/models/whisper_audio.cpp) (`VT_WHISPER_ENC_WARP=1`/`VT_WHISPER_ENC_EAGER=1` A/B) + [qwen3_vl_vision.cpp](../src/vllm/model_executor/models/qwen3_vl_vision.cpp) (vision tower default since §16; `VT_QWEN3VL_ATTN_WARP=1`/`VT_QWEN3VL_ATTN_EAGER=1` A/B) | **GPU GATE GREEN on dgx (2026-07-28, GB10 sm_121a, base `af1ed76b`):** CUDA `-Werror` 0-warn (cutlass-ON + FA2-ENABLED banners). `test_voxtral_e2e` **16/16** default-flash; flash/warp/eager token dumps md5-IDENTICAL (`89923566…`) ⇒ ZERO token flips; goldens md5 UNCHANGED (`voxtral_golden.json 8ab87b7e…`, `voxtral_neartie.json 937b9ad3…`, before==after). Proof-of-run nsys `AttentionDenseFlashKernel` 32 inst, ZERO `AttentionWarpKernel`/naive on encoder; RED confirmed (corrupt kernel → gate FAILS → restore → 16/16); `compute-sanitizer --tool memcheck` **0 errors**; 3 runs byte-identical. **A/B (same binary, `flock`, rep0 dropped):** attention **35.11 → 19.29 ms/layer (1.82×, NON-OVERLAPPING)**; encoder forward **~1834 → ~1375 ms (1.33×)**. **NOT at parity:** ~1.37 s vs vLLM ~43 ms TTFT (~32×, was ~44×) — the scalar warp-per-query recurrence is now serial-latency-bound over 1500 keys (L2 already served much of the redundant reads ⇒ 1.8× not 16×); gap-closer is a tensor-core MMA hd-64 non-causal FA2 instantiation (LARGE) + resident encoder weights (MEDIUM). **Vision tower (§16, 2026-07-28):** extended to the Qwen3-VL/27B tower (hd-72, 784 patches) — STRICT image/video e2e 32/32, bench flash-vs-warp 0/1,003,520 mismatch, nsys default 4B e2e `AttentionDenseFlashKernel` 24 inst/zero warp, RED 30/46→46/46, sanitizer 0; A/B warp 148.3→flash 142.3 ms = 1.04× (small — the vision attention at t=784 is serial-latency-bound not bandwidth-bound; the tower already BEATS vLLM at 0.57× eager) **2026-08-21 (`CLAIM-ATTN-RUNG-VISIBLE`, issue [#1544](https://github.com/mudler/vllm.cpp/issues/1544), spec [attention-rung-visibility.md](specs/attention-rung-visibility.md)): the advertised head_dim contract is now the LAUNCHABLE one, and the naive rung stops being a silent default.** The op stated `d <= 256` ([cuda_ops.cu](../src/vt/cuda/cuda_ops.cu) `LaunchAttentionDenseFlash`) while requesting `2*kFlashBc*d*sizeof(Tin)` bytes of DYNAMIC shared memory with no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`, so the driver's default 48 KiB cap made the real ceiling **192 bf16 / 96 f32** — Kimi (192 f32, 96 KB) and Qwen3.5 (256) would have taken a bare launch error from the `cudaGetLastError` at the bottom of the launcher, naming nothing they could do instead. The bound now lives in [ops.h](../include/vt/ops.h) as `AttentionDenseFlashSmemBytes` / `AttentionDenseFlashMaxHeadDim`, PURE host arithmetic so a box with no GPU can execute it, tied to the kernel by two `static_assert`s on `kFlashBc` and the register blocking; the launcher refuses above it naming `vt::AttentionDenseFast`, which uses NO shared memory and does serve those widths. NARROWING was chosen over `cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize)`: d=256 f32 wants 128 KiB, above the opt-in per-block cap of the consumer Blackwell parts gated here, so the opt-in would still leave the widest advertised width a lie AND cannot be verified without a device. Strictly additive for callers — the bound is INCLUSIVE, so d=192 bf16 lands exactly on 49152 and still launches. Mirrors `supports_head_size` / `get_supported_head_sizes` [backend.py:155-163](https://github.com/vllm-project/vllm/blob/555967922/vllm/v1/attention/backend.py#L155), consulted BEFORE dispatch rather than discovered by launching. Same change adds [check-attention-rung-consistency.py](../scripts/check-attention-rung-consistency.py) (preflight + CI): a model TU naming `vt::Attention` needs a `// VT-ATTN-NAIVE:` reason beside the call, so the six deliberate sites now say why and a new author gets a red instead of a silent ~500x. `kAttention` and every existing caller's numerics are UNTOUCHED by construction — the checker executes no model code and the head_dim guard only fires where the launch already failed. CPU-GATED: checker green, with unit + mutation cases in [test_check_attention_rung_consistency.py](../tests/scripts/test_check_attention_rung_consistency.py). NEITHER the site triple this cell used to quote (9 sites / 6 marked / 3 excused) NOR that suite's case count is recorded here, for one reason: a count of another file stored in this one reds on every row that legitimately adds or removes a `vt::Attention` call, which is the drift lock AGENTS.md `## Records` forbids and the same defect #1629 names inside that suite. Run the checker for the live triple and the suite for the live count, new head_dim contract cases in [test_ops_attention.cpp](../tests/vt/test_ops_attention.cpp). OWED [#1573](https://github.com/mudler/vllm.cpp/issues/1573): the on-device refusal case and its reachability mutation are PENDING a lease — `dgx:gpu0` was held by the developer, and the CPU cases pin the arithmetic, never that the launcher calls it. | [multimodal-speed](specs/multimodal-speed.md) §14 + §16 | `ACTIVE` | `CLAIM-MM-SPEED-AUDIO-ENC-KERNEL` + `CLAIM-MM-SPEED-QWEN-IMAGE` + `CLAIM-ATTN-RUNG-VISIBLE` |
 | `KERNEL-MOE-ROUTING` | Router top-k, align, permute/unpermute, combine, activation | core MoE sources `CMakeLists.txt:1135-1157`; M=1 decode parallelization mirrors `topk_softmax_kernels.cu:192-242,494-537` (moeTopK/topkGating) + `moe_align_sum_kernels.cu:147-185,295-324`; **grouped-topk (`noaux_tc`)** `fused_moe/router/grouped_topk_router.py:106-161` (`forward_native`; the fused `ops.grouped_topk` at `:28-70` is the same formula), upstream tests `tests/kernels/moe/test_grouped_topk.py`, `test_routing.py` | [cuda_moe.cu:349](../src/vt/cuda/cuda_moe.cu#L349); parallel router argmax [cuda_moe.cu:61](../src/vt/cuda/cuda_moe.cu#L61); parallel moe_align BlockScan [cuda_marlin_repack.cu:224](../src/vt/cuda/cuda_marlin_repack.cu#L224); **grouped-topk (MLA campaign W3)** — additive `MoeRouterTopKArgs` fields + optional `e_score_correction_bias` arg [ops.h](../include/vt/ops.h), CPU ref `MoeRouterGroupedTopKKernel` [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp) + CUDA `MoeRouterGroupedTopKKernel` [cuda_moe.cu](../src/vt/cuda/cuda_moe.cu). A SEPARATE kernel: `num_expert_group == 0` still dispatches the original, so the existing router is byte-identical by construction | [routing tests](../tests/vt/test_ops_moe.cpp#L57); byte-exact router+align parity [test_ops_moe_grouped.cpp:451](../tests/vt/test_ops_moe_grouped.cpp#L451); 35B 315/315 gate; **grouped-topk** [test_ops_moe_router_grouped.cpp](../tests/vt/test_ops_moe_router_grouped.cpp) at REAL DeepSeek-V3 dims (256 experts, n_group=8, topk_group=4, top_k=8, sigmoid, routed_scaling 2.5, WITH `e_score_correction_bias`) vs an INDEPENDENT sort-based transcription of the upstream formula, plus isolated cases for bias-selects/unbiased-weights, top-2-sum-vs-max group scoring, the group mask excluding the global argmax, and renorm-before-scaling; CPU-vs-CUDA ids EXACT + run-to-run bit-reproducible | [inventory](specs/kernel-family-inventory.md) | `ANCHOR-BACKFILL` | `CLAIM-MOE-DECODE-PARALLEL-1` |
 | `KERNEL-MOE-UNQUANTIZED` | Unquantized grouped/batched MoE GEMM | core MoE sources `CMakeLists.txt:1135-1157`; upstream `tests/kernels/moe/test_unquantized_backend_selection.py` | activation/combine subset [cuda_moe.cu:349](../src/vt/cuda/cuda_moe.cu#L349); grouped execution remains NVFP4-specialized | [MoE tests](../tests/vt/test_ops_moe.cpp#L193), [grouped tests](../tests/vt/test_ops_moe_grouped.cpp#L160) | [inventory](specs/kernel-family-inventory.md) | `PARTIAL` | - |
 | `KERNEL-MOE-QUANTIZED` | FP8/INT8/NVFP4/MXFP4 grouped MoE | CUTLASS/FP4 builds `CMakeLists.txt:865-1002`; NVFP4 oracle `fused_moe/oracle/nvfp4.py:38-276` | NVFP4 fallback [cuda_matmul_nvfp4.cu:761](../src/vt/cuda/cuda_matmul_nvfp4.cu#L761), Marlin [cuda_moe_marlin.cu:156](../src/vt/cuda/cuda_moe_marlin.cu#L156) | [NVFP4 grouped tests](../tests/vt/test_ops_moe_grouped.cpp#L160); 35B gate | [inventory](specs/kernel-family-inventory.md) | `PARTIAL` | - |
diff --git a/.agents/specs/attention-rung-visibility.md b/.agents/specs/attention-rung-visibility.md
index 1a7a94afd..1a24f9aa8 100644
--- a/.agents/specs/attention-rung-visibility.md
+++ b/.agents/specs/attention-rung-visibility.md
@@ -301,8 +301,21 @@ would be a regression rather than a repair.
   `scripts/check-attention-rung-consistency.py` ship unrepaired beside a suite
   that was repaired for #1629: `:58-61` gives the wrong reason for not widening
   the regex, `:93-96` attributes the exclusion of the fast rungs to the `\b`
-  rather than to the trailing `\(`, and `:252-255` denies an equality that holds
-  on this tree (9 sites, 6 marked, 3 excused). It cannot be fixed here, because
+  rather than to the trailing `\(`, and `:252-255` denies an equality that
+  holds ON EVERY GREEN RUN, for ANY allowlist, and always did -- which is a
+  larger claim than the one this entry used to make about one tree, and it is
+  the one #1631 needs. The comment's stated reason is that "a marked call inside
+  an allowlisted file counts in `marked`". True, and it does not separate the
+  quantities: such a call is counted in `marked` AND excluded from `excused`, so
+  it cancels on both sides. The only shape that separates `excused` from
+  `sites - marked` is an UNMARKED call in a NON-allowlisted file -- which is
+  exactly `drift_sites`, so `main` returns 1 at `:248` and the OK line never
+  prints. Enumerated rather than argued: over all 64 green configurations of
+  marked and unmarked calls across one allowlisted and one non-allowlisted file,
+  zero break the equality, and the one configuration that breaks it is not
+  green. So the printed `excused` is never anything but `sites - marked`, and
+  the comment justifying a separate computation is wrong wherever a reader can
+  see it. It cannot be fixed here, because
   changing what the gate accepts is what AGENTS.md `## Changing the rules or a
   checker` routes to its own row, spec and red-before evidence; attaching the
   correction to an unrelated semantic change is the alternative that section
@@ -310,6 +323,26 @@ would be a regression rather than a repair.
 
 ## Now
 
+**The allowlist is empty, and the three stems it carried are discharged (#1663,
+2026-08-22).** `## Scope`, the site table and `## Dependencies` above describe
+`ltx2`, `ltx2_device` and `muse_glimmer_vision` as rows in flight; they have
+landed. `47a918d8f` (#1579, issue #1545) routed Muse Glimmer's perception encoder
+to `vt::AttentionDenseFlash`, so it names `vt::Attention` nowhere; `90e8c3c85`
+(#1557, issue #1549) swapped the LTX-2.5 DiT device forward, and the two calls
+that remain in `ltx2.cpp` and `ltx2_device.cpp` are the host CPU-only arm and the
+`VLLM_LTX2_DIT_FLASH_ATTN=0` A/B arm, each now carrying its own marker.
+
+D4 predicted the checker would report those entries STALE and not fail, and it
+did, for the whole window from `90e8c3c85` to #1663. What D4 did not say is what
+that window COSTS: a listed stem excuses its entire translation unit, so deleting
+the live marker at `ltx2.cpp:959` left the checker at rc=0 (`7 carry a recorded
+reason, 1 unmarked and excused`) while the stem sat there, and reds at rc=1
+naming `ltx2.cpp:966` once it is gone. Both arms measured on `db648fb88` and
+restored byte-for-byte against a pre-taken sha256. The deferral D4 designs for is
+still right -- it keeps the removing row off this file -- but its cost is a real
+blind spot in the covered files, not only untidiness, and the next row that parks
+a stem should read it that way.
+
 The change is written, CPU-gated and through one fresh scoped review, whose
 findings are repaired here: the new checker registers its disabled creation-
 mutation stub in `check-pr-size.py` (measured 31 of 31 cases red under the stub);
diff --git a/scripts/attention-rung-allowlist.txt b/scripts/attention-rung-allowlist.txt
index 22041a1e5..bf1667713 100644
--- a/scripts/attention-rung-allowlist.txt
+++ b/scripts/attention-rung-allowlist.txt
@@ -28,9 +28,25 @@
 # that test too.
 
 # --- IN FLIGHT: the naive call is the defect, and another row is removing it ---
-muse_glimmer_vision   # 50 layers, H=16, head_dim=96, non-causal, sole path, no knob.
-                      # Issue #1545; the fix is vt::AttentionDenseFlash (24 KB of
-                      # shared memory at head_dim 96, inside the honest bound).
-ltx2                  # DiT self-attention; 47.84 s measured per forward (#1544).
-                      # LTX-2.5 routing row in flight.
-ltx2_device           # the device arm of the same DiT forward, same row.
+#
+# EMPTY, and that is the enforcement closed rather than the file retired. Every
+# stem this file ever held was deleted once its removing row landed:
+#
+#   muse_glimmer_vision  47a918d8f (#1579, issue #1545) routes the perception
+#                        encoder's sole path to vt::AttentionDenseFlash, so the
+#                        naive call is GONE rather than marked.
+#   ltx2, ltx2_device    90e8c3c85 (#1557, issue #1549) swaps the device DiT
+#                        forward to vt::AttentionDenseFlash. Both files still
+#                        NAME vt::Attention, and each surviving call now records
+#                        its own reason: ltx2.cpp's is the host arm, CPU-only by
+#                        construction, where kAttention and kAttentionDenseFlash
+#                        are the same registered function; ltx2_device.cpp's is
+#                        the VLLM_LTX2_DIT_FLASH_ATTN=0 arm of a same-binary A/B.
+#
+# Deleted here under #1663, by the preflight run the header above hands the
+# deletion to.
+#
+# An empty parking lot is the state to keep. Adding a stem is a review decision
+# and reds tests/scripts/test_check_attention_rung_consistency.py
+# ::ShippedTreeTests::test_allowlist_holds_only_the_in_flight_stems until the
+# expected set there is updated in the same change.
diff --git a/tests/scripts/test_check_attention_rung_consistency.py b/tests/scripts/test_check_attention_rung_consistency.py
index c4179d934..3e0f3b3e2 100755
--- a/tests/scripts/test_check_attention_rung_consistency.py
+++ b/tests/scripts/test_check_attention_rung_consistency.py
@@ -213,13 +213,53 @@ def test_the_six_deliberate_sites_carry_a_marker(self) -> None:
     def test_allowlist_holds_only_the_in_flight_stems(self) -> None:
         # It is not a parking lot. Growth is a review decision, and this pins the
         # set so growth is visible in a diff of this file.
+        #
+        # The set is EMPTY, which is the enforcement closed rather than the guard
+        # switched off: all three original stems were deleted once their removing
+        # rows landed (47a918d8f for muse_glimmer_vision, 90e8c3c85 for ltx2 and
+        # ltx2_device), and #1663 removed them here. An empty expected set still
+        # reds on the next silent append, which is the whole point of this case.
+        #
+        # It does NOT make the suite vacuous. drift_sites is what excuses a call,
+        # and with an empty allowlist it excuses nothing, so
+        # test_shipped_tree_is_green above now measures the tree on its markers
+        # alone. The stems it used to cover are re-asserted positively below.
         _, allowed = self.scan()
-        self.assertEqual(allowed, {"muse_glimmer_vision", "ltx2", "ltx2_device"})
+        self.assertEqual(allowed, set())
+
+    def test_the_formerly_allowlisted_stems_pass_on_their_own_merit(self) -> None:
+        # The obligation the allowlist deferred, now stated where a regression
+        # would be read: each of the three stems is green because its file earned
+        # it, never because a line in a text file excused it.
+        #
+        # ltx2 and ltx2_device still NAME vt::Attention -- a host CPU-only arm and
+        # the OFF arm of a same-binary A/B -- so they are asserted as MARKED.
+        # muse_glimmer_vision was routed to vt::AttentionDenseFlash outright, so
+        # it is asserted ABSENT from the scan. Asserting the same thing about all
+        # three would be false of one of them in either direction.
+        scanned, _ = self.scan()
+        for stem in ("ltx2", "ltx2_device"):
+            path = f"{MODELS}/{stem}.cpp"
+            self.assertIn(path, scanned, path)
+            self.assertTrue(all(marked for _, marked in scanned[path]), path)
+        self.assertNotIn(f"{MODELS}/muse_glimmer_vision.cpp", scanned)
 
     def test_every_allowlisted_stem_names_a_real_model_source(self) -> None:
         # Pins that every allowlisted stem names a model source that exists, so a
         # typo is reported at the typo.
         #
+        # DORMANT rather than dead now that the allowlist is empty: this iterates
+        # over nothing and asserts nothing on the shipped tree. That is the
+        # correct state for it -- it is a guard on a file that is currently
+        # empty, and it fires on the first TYPO added to it. Not on the first
+        # addition: both arms were measured, and they differ. Appending
+        # `zzz_bogus_model` reds this case and the pinning case together, while
+        # appending `whisper_audio` -- a real stem, so a real source file -- reds
+        # the pinning case and `test_deleting_a_marker_goes_red` and leaves THIS
+        # case green, because the stem it names exists. Deleting this case
+        # because it is quiet today would remove the typo report from exactly
+        # the edit that needs it.
+        #
         # Keyed on the FILE existing, never on scan membership. A stem stops having
         # a call site the moment its removing row lands -- that is the state the
         # allowlist is built to survive, which the checker's own