From 8bf6b57dbab3145d408913607387102ff48c5cbc Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 22 Aug 2026 07:03:02 +0000 Subject: [PATCH 1/3] fix(#1663): a stale attention-rung stem excused a DELETED marker in the file it named `47a918d8f` (#1579, issue #1545) and `90e8c3c85` (#1557, issue #1549) landed the routing the allowlist was parked for, and left their three stems behind. That is by design -- `scripts/check-attention-rung-consistency.py` reports a discharged stem as `STALE (not a failure)` and exits 0, so the removing row never has to edit this file and the deletion falls to the next preflight. This is that preflight. What the design did not say is what the window costs. A listed stem excuses its whole translation unit, not only the call that earned the entry, so every marker in a covered file is decorative while the stem sits there. Measured on `db648fb88`, deleting the `// VT-ATTN-NAIVE:` marker at `ltx2.cpp:958`: 3 stems listed: rc=0, "7 carry a recorded reason, 1 unmarked and excused" stems removed: rc=1, "src/vllm/model_executor/models/ltx2.cpp:966" Both arms restored byte-for-byte against a pre-taken sha256. So the entries were a live hole in the guard #1544 exists to be, open from `90e8c3c85` until here. The three earned their green independently, and asymmetrically: `muse_glimmer_vision` names `vt::Attention` nowhere, while `ltx2` and `ltx2_device` still name it at the host CPU-only arm -- where `kAttention` and `kAttentionDenseFlash` are the same registered function -- and at the `VLLM_LTX2_DIT_FLASH_ATTN=0` arm of a same-binary A/B, each carrying its own reason. The expected set in `test_allowlist_holds_only_the_in_flight_stems` moves to the empty set in the same change, which the allowlist header requires. Empty 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 rather than by their absence from a list. The header keeps the reason each stem left, so an empty parking lot does not read as an abandoned one. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/issue-index.md | 1 + .agents/specs/attention-rung-visibility.md | 26 ++++++++++++++-- scripts/attention-rung-allowlist.txt | 28 +++++++++++++---- .../test_check_attention_rung_consistency.py | 30 ++++++++++++++++++- 4 files changed, 76 insertions(+), 9 deletions(-) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index defaf1c0e..df4d5e59f 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -591,3 +591,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1646](https://github.com/mudler/vllm.cpp/issues/1646) | `SPEC-DFLASH2` | **`tests/tools/` held 351 cases across 20 suites and NO lane ran one of them.** Measured at `e100e64e1` on a clean worktree: `python3 -m unittest discover -s tests/tools -t . -p "test_*.py"` reports `Ran 351 tests in 21.067s / OK`, standard library only, zero skips, no GPU and no vLLM wheel — and no workflow, no CTest registration and no `scripts/agent-preflight.sh` line executed it. The only `unittest` invocations in `.github/workflows/` are four `tests.scripts.*` modules; preflight's `SUITES` loop runs `tests/scripts/$suite.py` only; a tree-wide grep for `tests.tools` outside the directory returns prose in `.agents/` and four unrelated path constants in `scripts/check-snapshot-pins.py`. Worse than untested: the suites are QUOTED AS EVIDENCE — `.agents/parity-ledger.md` carries "all tools 34/34" on five `SERVE-GATE-ONLINE` rows and `.agents/upstream-sync.md:38` records "34 of the 233 `tests/tools` cases" — so they read as gating in every document that cites them while being reachable only by an agent who typed the command. What they cover is not marginal: `test_oracle_pin.py` is the [#520](https://github.com/mudler/vllm.cpp/issues/520) oracle-identity assertion, `test_gpu_clock_state.py` is the [#543](https://github.com/mudler/vllm.cpp/issues/543) clock attribution every ratio rests on. FIXED IN FLOW by one preflight line and one CI step, DISCOVERED rather than enumerated because an enumerated list is a shared file every new suite must edit, which is the record-lock shape `AGENTS.md` §Records forbids. Found while wiring [#1562](https://github.com/mudler/vllm.cpp/issues/1562)'s refusal gate, which would otherwise have landed dead; not owned by that row's subject matter | bug | | [#1648](https://github.com/mudler/vllm.cpp/issues/1648) | `SPEC-DFLASH2` | **[#1646](https://github.com/mudler/vllm.cpp/issues/1646)'s central claim is FALSE: `tests/tools/` HAS a CTest registration and CI has been running it on every pull request.** #1646's row says the 351 cases across 20 suites had "no workflow, no CTest registration and no `scripts/agent-preflight.sh` line" executing them. `tests/CMakeLists.txt:12-16` registers the target `test_serve_low_tools` running exactly `python3 -m unittest discover -s tests/tools -p "test_*.py"` with `PYTHONPATH` set, and has since `e58858a91`; `CMakeLists.txt:2728` makes it live and `.github/workflows/ci.yml:1057` runs `ctest --test-dir build` inside `build-test-cpu` on every `pull_request`. Executing that command on the branch head reports `Ran 414 tests / OK` in 20.2 s. The error came from grepping `tests.tools`, the DOTTED module path, while CMake and the workflow spell it `tests/tools` -- a null grep proving the terms wrong rather than the thing absent. Three consequences: `.agents/parity-ledger.md`'s `all tools` citations and `.agents/upstream-sync.md:38` were citing a LIVE suite, not a dead one; the NARROW gap was real and its fix stands, because preflight genuinely ran none of them, and the CI step's remaining value is a Python-only verdict in ~21 s under its own name rather than behind a C++ build; and the DISCOVERY mechanism #1646 chose is correct and unaffected. TWO FURTHER ERRORS in the same immutable row, recorded here because it cannot be edited: "pinned at 3" pins nothing, since `hook_reasons` asserts `recorded_blocks <= propose_calls - skipped`, which holds for any residual at or above 0; and "all tools 34/34 on five rows" does not occur -- the ledger reads `35/35` on two rows and `164/164` on eleven, and 34 is `upstream-sync.md`'s count of cases that go RED under a released-wheel version shape, a different quantity from a pass count. The #1646 row is append-only and STANDS AS WRITTEN; this issue and `## Owed` O27 of [the DFlash2 spec](specs/dflash2-spec-decode.md) are the authority over it, as this campaign handled the #1538 row. Prose corrected in `.github/workflows/ci.yml` and `scripts/agent-preflight.sh` in the same change. Found by the fresh review of the SPEC-DFLASH2 speed harness ([#1562](https://github.com/mudler/vllm.cpp/issues/1562)) | bug | | [#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:958` 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 | diff --git a/.agents/specs/attention-rung-visibility.md b/.agents/specs/attention-rung-visibility.md index 1a7a94afd..286c85e30 100644 --- a/.agents/specs/attention-rung-visibility.md +++ b/.agents/specs/attention-rung-visibility.md @@ -301,8 +301,10 @@ 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 held + on the tree this spec was written against (9 sites, 6 marked, 3 excused) and + no longer does (8 sites, 8 marked, 0 excused, as of #1663) -- the comment is + still wrong, and it is now wrong about a different tree. 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 +312,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:958` 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..7800c19cd 100755 --- a/tests/scripts/test_check_attention_rung_consistency.py +++ b/tests/scripts/test_check_attention_rung_consistency.py @@ -213,8 +213,36 @@ 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 From f64effeca13c81855d55819742589d72ad5406f2 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 22 Aug 2026 07:47:28 +0000 Subject: [PATCH 2/3] fix(#1663): repair four fresh-review findings, one of which inverted a verdict F1, and it is the one that mattered. The `#1631` owed entry in `attention-rung-visibility.md` said the checker comment at `:252-255` "denies an equality that held ... and no longer does". Wrong, and wrong in the direction that misleads: if the equality no longer held, the comment denying it would be RIGHT. It still holds, and after this row it holds for a stronger reason -- with an empty allowlist, any green tree has `sites == marked` and `excused == 0` by construction, so the equality cannot fail. The same sentence also credited the 8/8/0 triple to #1663; `db648fb88` already read 8 sites and 8 marked before this row touched anything, because `47a918d8f` and `90e8c3c85` moved it. #1663 moved only the allowlisted-stem count, 3 to 0. F2. The marker is at `ltx2.cpp:959`, not `:958` -- `:958` is `a.causal = false;`. Corrected in the spec and in the index row, which is append-only and could not have been corrected after the merge. F3. The `KERNEL-ATTN-DENSE-FLASH` evidence cell still read "checker green (9 sites / 6 marked / 3 unmarked and excused by the 3 in-flight allowlisted stems)". Every number was false before this row and all three would be false again after it, which is the point: a count of one file stored in another reds on every row that legitimately adds or removes a call. The cell stops quoting the triple rather than restating it, since restating it rebuilds the drift lock AGENTS.md `## Records` forbids and #1629 already names in this checker's own test. The checker prints the live numbers; the suite pins what is invariant. F4. `test_every_allowlisted_stem_names_a_real_model_source` now iterates over an empty set. It is DORMANT, not dead, and the file says so: it guards a file that is currently empty and fires on the first thing added to it, which is measured rather than asserted -- appending one bogus stem reds it together with the pinning case. Gate after the repairs: checker rc=0, suite 35 of 35 OK, issue-index append-only rc=0 and `1 0` against `origin/main`. `test_cpu_x86_llamacpp_floor` reds at loadavg 31 with `NO_QUIET_WINDOW` and `busy=161%`, which is #618 and not reachable from a diff of three Markdown files and one Python comment. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/issue-index.md | 2 +- .agents/kernel-matrix.md | 2 +- .agents/specs/attention-rung-visibility.md | 13 ++++++++----- .../test_check_attention_rung_consistency.py | 8 ++++++++ 4 files changed, 18 insertions(+), 7 deletions(-) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index df4d5e59f..e04130450 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -591,4 +591,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1646](https://github.com/mudler/vllm.cpp/issues/1646) | `SPEC-DFLASH2` | **`tests/tools/` held 351 cases across 20 suites and NO lane ran one of them.** Measured at `e100e64e1` on a clean worktree: `python3 -m unittest discover -s tests/tools -t . -p "test_*.py"` reports `Ran 351 tests in 21.067s / OK`, standard library only, zero skips, no GPU and no vLLM wheel — and no workflow, no CTest registration and no `scripts/agent-preflight.sh` line executed it. The only `unittest` invocations in `.github/workflows/` are four `tests.scripts.*` modules; preflight's `SUITES` loop runs `tests/scripts/$suite.py` only; a tree-wide grep for `tests.tools` outside the directory returns prose in `.agents/` and four unrelated path constants in `scripts/check-snapshot-pins.py`. Worse than untested: the suites are QUOTED AS EVIDENCE — `.agents/parity-ledger.md` carries "all tools 34/34" on five `SERVE-GATE-ONLINE` rows and `.agents/upstream-sync.md:38` records "34 of the 233 `tests/tools` cases" — so they read as gating in every document that cites them while being reachable only by an agent who typed the command. What they cover is not marginal: `test_oracle_pin.py` is the [#520](https://github.com/mudler/vllm.cpp/issues/520) oracle-identity assertion, `test_gpu_clock_state.py` is the [#543](https://github.com/mudler/vllm.cpp/issues/543) clock attribution every ratio rests on. FIXED IN FLOW by one preflight line and one CI step, DISCOVERED rather than enumerated because an enumerated list is a shared file every new suite must edit, which is the record-lock shape `AGENTS.md` §Records forbids. Found while wiring [#1562](https://github.com/mudler/vllm.cpp/issues/1562)'s refusal gate, which would otherwise have landed dead; not owned by that row's subject matter | bug | | [#1648](https://github.com/mudler/vllm.cpp/issues/1648) | `SPEC-DFLASH2` | **[#1646](https://github.com/mudler/vllm.cpp/issues/1646)'s central claim is FALSE: `tests/tools/` HAS a CTest registration and CI has been running it on every pull request.** #1646's row says the 351 cases across 20 suites had "no workflow, no CTest registration and no `scripts/agent-preflight.sh` line" executing them. `tests/CMakeLists.txt:12-16` registers the target `test_serve_low_tools` running exactly `python3 -m unittest discover -s tests/tools -p "test_*.py"` with `PYTHONPATH` set, and has since `e58858a91`; `CMakeLists.txt:2728` makes it live and `.github/workflows/ci.yml:1057` runs `ctest --test-dir build` inside `build-test-cpu` on every `pull_request`. Executing that command on the branch head reports `Ran 414 tests / OK` in 20.2 s. The error came from grepping `tests.tools`, the DOTTED module path, while CMake and the workflow spell it `tests/tools` -- a null grep proving the terms wrong rather than the thing absent. Three consequences: `.agents/parity-ledger.md`'s `all tools` citations and `.agents/upstream-sync.md:38` were citing a LIVE suite, not a dead one; the NARROW gap was real and its fix stands, because preflight genuinely ran none of them, and the CI step's remaining value is a Python-only verdict in ~21 s under its own name rather than behind a C++ build; and the DISCOVERY mechanism #1646 chose is correct and unaffected. TWO FURTHER ERRORS in the same immutable row, recorded here because it cannot be edited: "pinned at 3" pins nothing, since `hook_reasons` asserts `recorded_blocks <= propose_calls - skipped`, which holds for any residual at or above 0; and "all tools 34/34 on five rows" does not occur -- the ledger reads `35/35` on two rows and `164/164` on eleven, and 34 is `upstream-sync.md`'s count of cases that go RED under a released-wheel version shape, a different quantity from a pass count. The #1646 row is append-only and STANDS AS WRITTEN; this issue and `## Owed` O27 of [the DFlash2 spec](specs/dflash2-spec-decode.md) are the authority over it, as this campaign handled the #1538 row. Prose corrected in `.github/workflows/ci.yml` and `scripts/agent-preflight.sh` in the same change. Found by the fresh review of the SPEC-DFLASH2 speed harness ([#1562](https://github.com/mudler/vllm.cpp/issues/1562)) | bug | | [#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:958` 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 | +| [#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 | diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 193ada943..550970a86 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. The counts this cell used to quote (9 sites / 6 marked / 3 excused) are NOT restated here, because a count of another file stored in this one reds on every row that legitimately adds or removes a `vt::Attention` call -- the drift lock AGENTS.md `## Records` forbids, and the same defect #1629 names in the checker's own test. The checker prints the live triple; `tests/scripts/test_check_attention_rung_consistency.py` pins what is invariant about it, 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-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 286c85e30..28b352977 100644 --- a/.agents/specs/attention-rung-visibility.md +++ b/.agents/specs/attention-rung-visibility.md @@ -301,10 +301,13 @@ 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 held - on the tree this spec was written against (9 sites, 6 marked, 3 excused) and - no longer does (8 sites, 8 marked, 0 excused, as of #1663) -- the comment is - still wrong, and it is now wrong about a different tree. It cannot be fixed here, because + rather than to the trailing `\(`, and `:252-255` denies an equality that + holds. It held when this spec was written (9 sites, 6 marked, 3 excused) and it + holds now for a STRONGER reason: `47a918d8f` and `90e8c3c85` moved the tree to + 8 sites and 8 marked, and #1663 then emptied the allowlist, so `excused` is 0 + BY CONSTRUCTION -- with nothing allowlisted, any green tree has + `sites == marked` and `excused == 0`, and the equality cannot fail. The comment + is still wrong, and it is now wrong about a tree where it cannot be right. 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 @@ -324,7 +327,7 @@ that remain in `ltx2.cpp` and `ltx2_device.cpp` are the host CPU-only arm and th 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:958` left the checker at rc=0 (`7 carry a recorded +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 diff --git a/tests/scripts/test_check_attention_rung_consistency.py b/tests/scripts/test_check_attention_rung_consistency.py index 7800c19cd..38fcdcaa7 100755 --- a/tests/scripts/test_check_attention_rung_consistency.py +++ b/tests/scripts/test_check_attention_rung_consistency.py @@ -248,6 +248,14 @@ 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 the moment anything is added, which is measured + # rather than assumed (appending one bogus stem reds this case and the + # pinning case together). Deleting it 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 From 7865f9280d0b420e747c7f2ff0affe337e68bc9b Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 22 Aug 2026 07:59:28 +0000 Subject: [PATCH 3/3] fix(#1663): the dormant case fires on a TYPO, not on an addition, and the equality is wrong on every green run Three repairs from the scoped re-review of `f64effeca`, two of which correct claims this branch itself made. The F4 comment overstated by one clause. It said the dormant `test_every_allowlisted_stem_names_a_real_model_source` "fires the moment anything is added", in a comment whose whole subject is that the claim is measured. Both arms are measured now and they differ: appending `zzz_bogus_model` reds that case together with the pinning case, while appending `whisper_audio` -- a real stem with a real source file -- reds the pinning case and `test_deleting_a_marker_goes_red` and leaves the dormant case GREEN. It fires on the first TYPO, which is what it was built to report, and the comment now says that. The parenthetical was already correct; the clause a skimmer takes away was not. The F1 entry UNDERSTATED, which is the direction that costs #1631 its evidence. It said the checker comment at `:252-255` is wrong about a tree where it cannot be right. The comment is wrong on EVERY green run, for ANY allowlist, and always was. Its stated reason -- "a marked call inside an allowlisted file counts in `marked`" -- is true and does not separate the quantities, because such a call is counted in `marked` AND excluded from `excused` and 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: 64 green configurations of marked and unmarked calls across one allowlisted and one non-allowlisted file, zero break the equality, and the single configuration that does break it is not green. F3's replacement text named the suite twice and spliced two clauses into something that read like a merge artefact. The cell now gives one reason for withholding two counts -- the site triple and the suite's case count -- instead of giving that reason twice in different words. Focused gate after the repairs: checker rc=0 at `8 sites / 8 marked / 0 excused by 0 allowlisted`, suite 35 of 35 OK, `check-issue-index-append-only.py` rc=0. Both mutation arms restored byte-for-byte against a pre-taken sha256. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/kernel-matrix.md | 2 +- .agents/specs/attention-rung-visibility.md | 20 +++++++++++++------ .../test_check_attention_rung_consistency.py | 12 +++++++---- 3 files changed, 23 insertions(+), 11 deletions(-) diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 550970a86..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. The counts this cell used to quote (9 sites / 6 marked / 3 excused) are NOT restated here, because a count of another file stored in this one reds on every row that legitimately adds or removes a `vt::Attention` call -- the drift lock AGENTS.md `## Records` forbids, and the same defect #1629 names in the checker's own test. The checker prints the live triple; `tests/scripts/test_check_attention_rung_consistency.py` pins what is invariant about it, 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 28b352977..1a24f9aa8 100644 --- a/.agents/specs/attention-rung-visibility.md +++ b/.agents/specs/attention-rung-visibility.md @@ -302,12 +302,20 @@ would be a regression rather than a repair. 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. It held when this spec was written (9 sites, 6 marked, 3 excused) and it - holds now for a STRONGER reason: `47a918d8f` and `90e8c3c85` moved the tree to - 8 sites and 8 marked, and #1663 then emptied the allowlist, so `excused` is 0 - BY CONSTRUCTION -- with nothing allowlisted, any green tree has - `sites == marked` and `excused == 0`, and the equality cannot fail. The comment - is still wrong, and it is now wrong about a tree where it cannot be right. It cannot be fixed here, because + 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 diff --git a/tests/scripts/test_check_attention_rung_consistency.py b/tests/scripts/test_check_attention_rung_consistency.py index 38fcdcaa7..3e0f3b3e2 100755 --- a/tests/scripts/test_check_attention_rung_consistency.py +++ b/tests/scripts/test_check_attention_rung_consistency.py @@ -251,10 +251,14 @@ def test_every_allowlisted_stem_names_a_real_model_source(self) -> None: # 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 the moment anything is added, which is measured - # rather than assumed (appending one bogus stem reds this case and the - # pinning case together). Deleting it because it is quiet today would - # remove the typo report from exactly the edit that needs it. + # 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