fix(SPEC-DFLASH2): build-newest-gcc has been red since W5, because the GGUF test bypassed the process_id seam (#1575) - #1581
Merged
Conversation
…e GGUF test bypassed the process_id seam (#1575) `build-newest-gcc` has been RED on `main` since `5702d8f83`, and the finding is not a missing include. It is a recurrence of the exact class `tests/support/process_id.h` was created to close. `tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547` calls `::getpid()` while including nothing that declares it. The lane builds in a `gcc:16` container, where libstdc++ no longer supplies the transitive `<unistd.h>`, so the file does not COMPILE: error: '::getpid' has not been declared; did you mean 'getpt'? Every local toolchain here is gcc 13 and still supplies that include, which is why the lane is the only place it shows. The seam predicted this in its own header: the spelling "was fixed once in three files and came back in five more, because each new loader test copies the temp-directory helper from the last one". W5's GGUF test copied the temp-directory idiom and the direct call with it. The convention is otherwise unanimous -- every other file in the tree naming `getpid` either includes `<unistd.h>` or uses the helper, and the sibling `tests/vllm/entrypoints/test_dspark_draft_routing.cpp` uses the helper AND carries a comment warning not to reintroduce this class. So the repair is the seam, not `<unistd.h>`. The POSIX spelling would build this lane green and re-copy the idiom the helper exists to centralise, and it does not compile on MSVC at all. ATTRIBUTED, not inferred: `build-newest-gcc` is absent from the baseline lane's failed list at `92406c620` and present at `5702d8f83`, the commit that added the file. MEASURED red-first in the lane's own container, at base `947e5f648`, file sha256 `83bba319...`: - unmodified, `gcc:16`, `-fsyntax-only`: RED_RC=1, error byte-identical to CI - with the seam applied: GREEN_RC=0, compile_err=0 - local gcc 13, same command: GCC13_RC=0, compile_err=0 `git diff --stat` was read on the green arm as well, so a mutation that silently failed to apply could not read as a pass. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
…tes skipped while the branch was behind (#1575) `scripts/agent-preflight.sh --fail-on-skip` SKIPPED `commit-trailers` and `commit-style` on this branch, because `origin/main` had moved to `4c193bd55` and was no longer an ancestor of HEAD. The skip reason says it plainly -- "Neither gate reported anything about this tree" -- so the green beside them was an absence of measurement rather than a pass. Merging `origin/main` at 5d548d0 makes both gates RUN. Nothing here conflicts with the fix: the two incoming commits touch the lifecycle doc-checkpoint gate and an MTP GGUF test, and this branch touches a DFlash2 test file, the DFlash2 spec header and one appended issue-index row. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
This was referenced Aug 21, 2026
Collaborator
Author
Fresh review — verdict PASSDispatched by the landing session into its own worktree, against the immutable Verdict: PASS. Four findings, none blocking, none requiring a repair before The mutation, in the lane's own containerThe claimed guarantee is "this file compiles under gcc 16", and it was mutated
The mutant's entire stderr is byte-identical to the CI error this body quotes: Both of this repo's standing mutation traps are cleared explicitly. It did not What the reviewer re-derived rather than accepted
Findings
Full gate
The one failure is The full suite was cited, not run — the box is memory-tight and Standing concern the reviewer asked be carried forwardDo not read a merge of this PR as "main is green". The newest baseline FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true |
This was referenced Aug 21, 2026
localai-bot
added a commit
that referenced
this pull request
Aug 21, 2026
…edundant (#1617) (#1618) `d27639e71` added `#include <unistd.h>` to test_qwen3_dflash2_gguf.cpp to declare a `::getpid()` call that `13548db8f` had already removed, three commits earlier, by routing the same need through `vllm_test::ProcessId()` and `support/process_id.h`. The change was authored against the pre-#1581 tree, where it was the correct fix, and landed after the call it supported was gone. The file now carries no `::getpid` call at all, so the include declares nothing it uses. It is also the one spelling `tests/support/process_id.h` exists to prevent: MSVC ships no `<unistd.h>`, so an unguarded include does not merely warn on Windows, it fails to compile, which is the #603 shape of the #503 class. It did not surface as a new red because both `windows-msvc-*` lanes are baseline-red and never run on `main`. Verified by compiling the translation unit with the project's own flags, including -Werror, after the removal: rc 0, object produced. The file still includes `support/process_id.h` and still calls `vllm_test::ProcessId()`. Found while merging external contributor pull requests and fixed in the same flow; it is not that work's defect. Closes #1617 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
…use TWO other changes landed it first and the merge kept all three `origin/main` reached `13548db8f` and the `::getpid()` defect this branch also repaired is now fixed there **twice**. `d27639e71` (#1498, a Tenstorrent decode fix) added `#include <unistd.h>` to `tests/vllm/models/test_qwen3_dflash2_gguf.cpp`, and `13548db8f` (#1581) routed the call through the `tests/support/process_id.h` seam. Three sessions repaired one defect under two issue numbers, #1565 and #1575. **The merge produced a file with `<unistd.h>` twice and no conflict.** This branch's include sits in a different part of the include block from `main`'s, so git auto-merged both and nothing warned. That is the silent shape: a conflict is visible and a clean automerge of two repairs for one defect is not. So this branch takes `main`'s version of that file **whole**. Its only change to it was `118c8adae`'s include, verified by `git log e2a9e03..HEAD -- <file>` returning that one commit, and the file is now byte-identical to `origin/main` (`git diff origin/main -- <file>` is empty). The bare `<unistd.h>` is exactly the copy-forward idiom the seam exists to centralise, and #1581's review measured the seam repair in the lane's own `gcc:16` container, so the seam is the version that should survive. Nothing is lost: `build-newest-gcc` is repaired on `main` and this branch inherits it. `Closes #1565` is therefore no longer this branch's to claim and comes out of the PR body in the same step. `.agents/issue-index.md` merged as a pure union-append, checked by multiset against both parents: 547 and 547 lines in, 555 out, no parent line dropped, altered or invented. `tests/vllm/multimodal/test_ltx2_video.cpp` is untouched by this merge. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
…he redundant include this branch had also dropped `main` moved twice while this branch was being reconciled. `13548db8f` (#1581) landed the `tests/support/process_id.h` seam and `369a0c2aa` (#1618) then dropped the bare `#include <unistd.h>` that #1498 had added for the same defect, so the file is back to one repair. This branch already held `main`'s version of it whole, and `git diff origin/main -- tests/vllm/models/test_qwen3_dflash2_gguf.cpp` is still empty after this merge. The merge is taken so `check-commit-style` and `check-commit-trailers` can RUN: both refuse a range whose base is not an ancestor of its head, and a moved `origin/main` turns them into SKIPs rather than passes. `.agents/issue-index.md` was the only file both sides touched and merged as a pure union-append, checked by multiset against both parents -- 555 and 549 lines in, 557 out, no parent line dropped, altered or invented. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
…h's CI can report on more than one broken lane `build-newest-gcc` failed at the BUILD step on this branch's last run with `test_qwen3_dflash2_gguf.cpp:547: error: '::getpid' has not been declared`, which is the `main` red #1581 repaired at `13548db8f`. A lane that does not compile reports nothing about this change, so the merge is what makes its verdict readable rather than a convenience. `.agents/issue-index.md` was the only file both sides touched and it merged as a pure union-append, checked by multiset against both parents: 532 and 548 lines in, 550 out, with no parent line dropped, altered or invented. Nothing else conflicted. `build-test-cpu` on the previous run also failed `456 - test_cpu_threadpool` on `CHECK( ratio < 100.0 )` at `128.28` -- the empty-op dispatch timing guard, which PASSED in 0.34 s on the `main` baseline run and which nothing in this diff can reach. That was not accepted as inherited on one observation; this run re-rolls it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
…e gcc lane can compile and report on this change `build-newest-gcc` failed at the BUILD step on this branch's last run with `test_qwen3_dflash2_gguf.cpp:547: error: '::getpid' has not been declared`, which is the `main` red #1581 repaired at `13548db8f`. A lane that does not compile reports NOTHING about this change, so the merge is what makes its verdict readable rather than a convenience. `.agents/issue-index.md` was the only file both sides touched and it merged as a pure union-append, checked by multiset against both parents: 550 and 549 lines in, 551 out, no parent line dropped, altered or invented. `build-test-cpu` on the previous run also failed `456 - test_cpu_threadpool` on `CHECK( ratio < 100.0 )` at `128.28` -- the empty-op dispatch timing guard, which PASSED in 0.34 s on the `main` baseline run and which nothing in this diff can reach. One observation is not an inheritance proof, so this run re-rolls it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
…that will land Brings the branch from `503e459005` to `2e7f3bee7`, seventeen commits, so the gate runs on the tree that will land rather than on the one the wave opened. Three of them change what the gate reads: `13548db8f` (`SPEC-DFLASH2`, #1581) and `369a0c2aa` (`SPEC-DFLASH2`, #1617) rewrite this row's own `tests/vllm/models/test_qwen3_dflash2_gguf.cpp`, and `52e328789` (`GATE-DOC-CHECKPOINT-STATES`, #1621) rewrites `scripts/check-doc-checkpoint.py`, the gate that reads this row's `docs/STATUS.md` lifecycle move. Gating the pre-merge tree would gate none of them. ONE conflict, `.agents/benchmark-record.md`, and it is a genuine both-sides-appended EOF conflict rather than a contended entry. Verified before resolving, not assumed: `git diff --numstat` from the merge base reads `79 0` for main and `243 0` for this row -- zero deleted lines on either side -- and this row's file is byte-for-byte the merge base's 25990 lines followed by 243 new ones, proven with `cmp` rather than by reading the diff. Main's 79 lines are one section inserted at 25566 and one appended at EOF, also with no line removed. Resolved as main's complete 26069-line file followed by this row's 243-line append: 26312 lines, with main re-verified by `cmp` as a byte-exact prefix and the appended block re-verified as byte-identical to this row's. `.agents/issue-index.md` NEEDED A MANUAL RESOLUTION EVEN THOUGH IT DID NOT CONFLICT. The `merge=union` driver produced a file that silently DROPPED main's last row, `#1502` (`VT-REFTIER-HOST-ADDRESSABLE`, landed by `c477f6f12`): 514 rows at the merge base, 534 on main, 520 here, and the automatic result held 537 where 540 is correct. Neither side had edited or deleted anything -- both `git diff --numstat` readings are `N 0` -- so this is the driver losing a row at the end of the file, not a contended record. Re-derived by COUNTING the merged rows rather than by arithmetic: main's complete 534-row file plus this row's 6 new rows appended at the end, 540 rows, main byte-exact as a prefix, every row of both parents present, no duplicate issue number, and `git diff --unified=0` against `2e7f3bee7` removes zero lines. The other seven files both sides touch were three-way merged and each was then verified by reading the merged bytes: for every one of them the merged file differs from `2e7f3bee7` by exactly this row's edit and differs from `635e7b5f9` by exactly main's edit. No key was contended. `.agents/engine-matrix.md` is the case that mattered, and the keys are disjoint: main rewrote `ENG-MM-INPUT-PIPELINE`, this row rewrote `SPEC-DFLASH2`. Every summary counter was re-derived by counting the merged section rows with the repository's own parser (`scripts/check-agent-record.py` `parse_claim_rows`), never by adjusting either side's number: 169 rows, ANCHOR-BACKFILL 35, PARTIAL 18, SPIKE 4, READY 13, ACTIVE 40, GATING 8, DONE 12, INVENTORIED 38, BLOCKED 1, which is what the `**Total**` row states. BLOCKED has no column, which is why the row count exceeds the column sum by one. `docs/BENCHMARKS.md`, `docs/FEATURES.md` and `docs/STATUS.md` each had both sides edit a row, and in each the two keys are different. Rows: SPEC-DFLASH2-W6, GATE-QWEN38-27B-FP8-BLOCK, VT-REFTIER-HOST-ADDRESSABLE, ENV-LEASE-GPU-CAPABILITY, SERVE-C-ABI, BACKEND-ROCM, GATE-DOC-CHECKPOINT-STATES, BACKEND-TENSTORRENT-HOST-FREE-FORWARD, FIX-GCC15-ARRAY-BOUNDS, ENG-MM-INPUT-PIPELINE FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [Claude Code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
… tree that will land Brings the branch from `f4ccabbb4` to `2e7f3bee7`, thirteen commits, so the gate runs on the tree that will land rather than on the one the wave opened. Two of them change what this row's gate reads: `13548db8f` (`SPEC-DFLASH2`, #1575, #1581) moves the GGUF DFlash2 test onto the `process_id` seam and `369a0c2aa` (`SPEC-DFLASH2`, #1617) trims its include, and both sit in the speculative-decoding suite this row extends. `b80659276` (`BACKEND-ROCM`, #1046) and `49c64bbc8` (`BACKEND-TENSTORRENT-HOST-FREE-FORWARD`, #1514) each add CTest entries, so the pre-merge tree runs a smaller suite than the one that lands. NO CONFLICT, and the clean merge was still verified rather than trusted. Five files are touched by both sides, and for each the merged bytes were read: the merged file differs from `2e7f3bee7` by exactly this row's edit and differs from `e1cd772f1` by exactly main's edit. `.agents/issue-index.md` is append-only under `merge=union` and was re-derived by COUNTING the merged rows, not by arithmetic on either side's number: 520 rows at the merge base, 534 on main, 525 here, 539 merged, which is main's 534 plus this row's 5. Zero deletions, zero modified rows, no duplicate issue number, every row of both parents present byte-for-byte, and `git diff --unified=0` against `2e7f3bee7` removes zero lines. The count was checked because the same driver silently DROPPED main's last row on the sibling `row/SPEC-DFLASH2-W6` merge taken from the same main SHA; it did not do so here, and nothing but a row count would have said which. `docs/STATUS.md` had both sides edit a row and the keys are disjoint: main rewrote `Tenstorrent Blackhole`, this row rewrote `Speculative decoding`. `docs/BENCHMARKS.md` is main's file plus this row's single appended `SPEC-DRAFTER-CHAIN` line, with main's rewritten host-free-decode row intact. `docs/USAGE.md` and `tests/CMakeLists.txt` interleave additions only. `.agents/engine-matrix.md` is the one keyed record main did not touch in this range, so this row's version stands unopposed. Its summary counters were still re-derived by counting the merged section rows with the repository's own parser (`scripts/check-agent-record.py` `parse_claim_rows`): 169 rows, ANCHOR-BACKFILL 35, PARTIAL 18, SPIKE 4, READY 12, ACTIVE 41, GATING 8, DONE 12, INVENTORIED 38, BLOCKED 1, which is what the `**Total**` row states. BLOCKED has no column, which is why the row count exceeds the column sum by one. Rows: SPEC-DRAFTER-CHAIN-W1, GATE-QWEN38-27B-FP8-BLOCK, VT-REFTIER-HOST-ADDRESSABLE, ENV-LEASE-GPU-CAPABILITY, SERVE-C-ABI, BACKEND-ROCM, GATE-DOC-CHECKPOINT-STATES, SPEC-DFLASH2, BACKEND-TENSTORRENT-HOST-FREE-FORWARD FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [Claude Code]
localai-bot
pushed a commit
that referenced
this pull request
Aug 21, 2026
… the trailer gates stop skipping `origin/main` was 38 commits ahead at `5453e571d`, which made `scripts/agent-preflight.sh --fail-on-skip` report `commit-trailers` and `commit-style` as SKIP: both gates refuse to run when `origin/main` is not an ancestor of `HEAD`, so neither had said anything at all about this tree. A skip is not a pass, so the sync is the only way to get a verdict. `git merge-tree --write-tree` was clean before the merge and the merge itself reported no conflicted path. The one keyed record both sides touch is `.agents/issue-index.md`, and it was checked by key rather than trusted: 520 rows at the merge base, 522 on this branch (the two rows this row appended), 566 on `main` (46 appended there), 568 merged. `sort | uniq -d` reports no duplicated row, and the #1421 row this row owns appears exactly once, at line 478, byte-identical to the version this branch committed. That is the append-only shape the union merge is allowed to take, not an accepted three-way merge of a keyed record. This sync also clears one inherited red rather than one this row caused. `build-newest-gcc` was failing on `main` and is fixed there by #1581, which moved the callers onto `vllm_test::ProcessId()`; #1565 is closed. That red on this pull request was stale CI from before the fix and comes back green from this merge. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode]
localai-bot
pushed a commit
that referenced
this pull request
Aug 22, 2026
…e this branch reasons about (#1074) `origin/main` moved to `5453e571dd412e63de53bd6ae66b5a94cd7df1e3` while this branch's CI was running, so the trailer gates would SKIP again and `build-newest-gcc` was still running against a base that predates its own fix. A SKIP reports nothing about this tree. Three of the 32 matter to this pull request beyond the merge. #1577, #1609 and #1621 all change `scripts/check-doc-checkpoint.py`'s record states, and this branch's description asserted that the gate's `STATES` tuple could see neither `INVENTORIED` nor `PARTIAL`. That claim is re-derived against the merged checker rather than carried, because a stale premise about a gate is the defect this branch already had once. #1581 landed the `process_id.h` seam that had reddened `build-newest-gcc` since W5, so that job's earlier red on this branch was stale CI rather than a finding. Four files this pull request edits also moved on `main`: `.agents/issue-index.md`, `docs/BENCHMARKS.md`, `docs/FEATURES.md` and `docs/STATUS.md`. The append-only and public-document gates are rerun on the merge result rather than assumed from a clean auto-merge. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot
added a commit
that referenced
this pull request
Aug 22, 2026
…ow says why, and AttentionDenseFlash advertises the head_dim it can launch (#1544) (#1578) Two additive changes from #1544's `## Owed`. Neither moves a single existing caller's numerics, and that constraint shaped the whole design. ## The naive rung stops being a silent default `vt::Attention` resolves `OpId::kAttention` straight to the correctness-grade kernel, and nothing in the tree ever routes it up: the rung is whichever C++ function name the author typed. That is deliberate for six of the nine call sites, and invisible to everyone else, which is how one LTX-2.5 DiT forward came to cost 47.84 s. A token gate cannot see the difference by construction — every rung is bit-identical or inside the bf16 envelope, so the goldens pass either way. A selector that auto-routes was rejected, and not on taste. Three of the six sites are reference arms a gate compares against (`nemotron_h.cpp`, `nemotron_h_device.cpp`, `qwen3_5.cpp`), two are the `VT_*_EAGER` rungs of a same-binary A/B (`whisper_audio.cpp`, `qwen3_vl_vision.cpp`), and one is a measured-negative device path behind `VT_KIMI_DEVICE_MLA`. Rerouting any of them changes what the reference computes, which deletes the comparison the gate performs rather than fixing anything — the "widen the assertion until the gate passes" failure AGENTS.md names. `kAttention` and `vt::Attention` are untouched here. `scripts/check-attention-rung-consistency.py` requires the CHOICE to be recorded instead: a `// VT-ATTN-NAIVE:` reason on the call line or within 20 lines above it. The six deliberate sites now carry one, and an author who never heard of the fast rungs gets a red instead of a silent 500x. The scan runs over `checker_text.normalize_source`, so a commented-out, `#if 0`-ed or `if (false)`-ed call is a deletion to it exactly as it is to nvcc, and the reported `file:line` still describes the original file. The record is per-site and in-file, so an ordinary change writes no shared record at all. `scripts/attention-rung-allowlist.txt` holds only the three stems whose naive call another row is currently deleting — `muse_glimmer_vision` (#1545) and the two LTX-2.5 files — because editing the very lines those changes replace would conflict for no gain. A stale entry there is reported and is not fatal, so the removing row owes this file nothing. ## `AttentionDenseFlash` advertises the head_dim it can launch It claimed `head_dim <= 256` while asking the driver for `2*kFlashBc*d*sizeof(Tin)` bytes of dynamic shared memory, with no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`. The default 48 KiB cap made the real ceiling 192 in bf16 and 96 in f32, so Kimi at 192 f32 or Qwen3.5 at 256 would have received a bare launch error from the `cudaGetLastError` at the bottom of the launcher, naming nothing they could do instead. The bound now lives in `include/vt/ops.h` as `AttentionDenseFlashSmemBytes` and `AttentionDenseFlashMaxHeadDim` — pure host arithmetic, so a box with no GPU can execute it — tied to the kernel by two `static_assert`s on the tile width and the register blocking. The launcher refuses above it naming `vt::AttentionDenseFast`, which uses no shared memory and does serve those widths. Narrowing beats opting in to a larger cap here, and that is now a measurement rather than a preference. **GB10's queried opt-in ceiling is 101,376 bytes** (measured during #1557's review), while head_dim 256 in f32 wants 131,072. So `cudaFuncSetAttribute` cannot make the widest advertised width true on the part this project gates on — the raise buys nothing at the width that motivated it, and a caller there would have gone on falling back silently without ever launching. The bound is INCLUSIVE, which matters in one direction: head_dim 192 in bf16 lands exactly on 49152 and launches today, so an exclusive bound would refuse work that currently runs. Opting in stays available later as a widening for bf16 above 192, owned by nobody today. This mirrors vLLM's own polarity rather than inventing one: `vllm/model_executor/models/vision.py:99` selects an encoder backend by shape, and `vllm/v1/attention/backend.py:155-163` consults `supports_head_size` BEFORE dispatch instead of discovering the domain by launching. Both read at the pinned oracle `555967922`. ## Evidence RED first, on the unmodified tree: the checker reported all six deliberate sites at the exact lines #1544 names (`kimi_linear_device.cpp:598`, `nemotron_h.cpp:671`, `nemotron_h_device.cpp:330`, `qwen3_5.cpp:5279`, `qwen3_vl_vision.cpp:527`, `whisper_audio.cpp:324`) and correctly excluded the three allowlisted ones. GREEN after the markers: 9 sites, 6 marked, 3 allowlisted. `tests/scripts/test_check_attention_rung_consistency.py` 34/34, including six mutations that must go RED — a new unmarked model, a new unmarked call in a HEADER, a deleted marker, a second unmarked call inside an already-marked file, a stub reason, and a widened regex that would swallow the fast rungs. It also pins that the scanned population is not empty, which is the guard against the way a structural checker usually goes green: by matching nothing at all, and that every allowlisted stem names a model source that exists, which is what catches a typo. `tests/vt/test_ops_attention.cpp` gains the head_dim contract cases: the tile arithmetic at both element sizes, both honest bounds, that 256 is outside both, and the inclusive edge in both directions. 11 cases / 39 assertions, SUCCESS. MUTATED, because a green suite over new arithmetic proves only that the arithmetic agrees with itself. Making `AttentionDenseFlashMaxHeadDim` return `kAttentionDenseMaxHeadDim` — the exact contract this change repairs — turns the new case RED on 6 of its assertions, each printing the wrong value it now carries (`256 == 192`, `256 == 96`, `65536 <= 49152`, `131072 <= 49152`), so the mutation demonstrably applied and demonstrably compiled. `include/vt/ops.h` restored and verified by sha256 against its pre-mutation snapshot; rebuilt; 39/39 green again. Proof that no caller's numerics moved: the checker executes no model code; the head_dim guard fires only where the launch already failed; no marker changes a statement; and `git diff` touches no kernel arithmetic, no dtype and no default. ## What the fresh review changed Six findings, repaired here. None of them moves a kernel's arithmetic either. `pr-size` was RED, and this branch caused it. A checker created inside the range has no BASE version for the red-before half of the evidence run, so it has to register the disabled stub its own suite must reject; about twenty checkers do, and this one did not, so the gate could not classify the change at all. Measured rather than asserted: under the stub every case goes red — re-measured after the repairs below, `FAILED (errors=34)` — because the suite loads the checker as a module and every case calls into it. `agent-preflight.sh` does not run `check-pr-size`, which is why a local green said nothing about it. The second `static_assert` beside the kernel was a tautology. It read `8 * 32 == kAttentionDenseMaxHeadDim` while the real `kMaxPerLane` was a function-local `constexpr` inside the kernel body, invisible at file scope, so setting that local to 4 — precisely the drift the message claims to catch — left the assert reading `256 == 256`. The register blocking is now `kFlashMaxPerLane` at file scope; the kernel's register arrays and unrolled loops read it, and so does the assert. The same mutation now reads `128 == 256` and fails to compile. There is no nvcc on this box, so the tie was measured by extracting that constant block from `cuda_ops.cu` VERBATIM and compiling it against the shipped `include/vt/ops.h` under `g++ -fsyntax-only`: clean before, `static assertion failed` after, `cuda_ops.cu` restored and verified by sha256. The first assert (`kFlashBc == kAttentionDenseFlashTileCols`) was already a real tie and is untouched. Two comments claimed more than the code delivers. The launcher said its guard and its shared-memory request came from "the SAME function … cannot disagree"; they are two functions, and `AttentionDenseFlashMaxHeadDim` re-derives the division instead of inverting `AttentionDenseFlashSmemBytes`. The guarantee holds and is tested: mutating the `2 *` in `SmemBytes` to `3 *` reds 9 assertions of the shipped contract case, both inclusive-edge checks among them, while `MaxHeadDim(2) == 192` stays green — which is the re-derivation made visible. The comment now describes that. The `AttentionDenseFa2` fall-through comment promised "the best available kernel for their shape rather than a hard refusal", which stopped being true for an over-cap head_dim the moment this branch added the refusal; it now names the domain and records that every caller today is far inside it (max head_dim 80). The checker claimed "the population is what makes a green meaningful" and named no limits. Four spellings reach the same kernel undetected — a `using` declaration, a namespace alias, a `#define`, and a call through `&vt::Attention` — each verified during review to leave the checker green with a live unmarked call. None exists in this tree, and widening the regex would make every fast rung a site, which is D1's rejected failure mode again; closing it needs a compiler-side population, not a longer pattern. The docstring and spec D6 state the bound, so a green reads as "no unmarked `vt::Attention(` call" and never as "no model is naive". The OK line reported total and marked sites but never the number a reader needs: sites carrying no reason that pass only because their stem is allowlisted. It is not `sites - marked`, since a marked call inside an allowlisted file counts in `marked`. Two cases now pin the line; dropping the count from it reds them. Two records were wrong. `scripts/attention-rung-allowlist.txt` told a removing row to delete its stem without saying that `test_allowlist_holds_only_the_in_flight_stems` pins the set in another file and reds on the deletion; the allowlist header, the checker docstring and spec D7 now say so. The kernel-matrix cell stored this suite's case count — a measurement of one file inside another, which AGENTS.md names as a drift lock — so the count is gone rather than corrected. ## Two drift locks in the new suite, both repaired here Found while landing this change against #1579, by checking the interaction instead of assuming the two pull requests were independent. Each was green on its own; `main` went red only once both landed, which is why nothing on either branch caught it. **The first was the population floor.** `test_the_population_is_not_empty` asserted the scanned population was `>= 9` against a tree of EXACTLY 9 sites. Its own name says "is not empty" and its assertion pinned a count: the name was right. A raw total is a measurement of the model tree stored in a test file, which AGENTS.md `## Records` forbids, and it reds on any row that legitimately REMOVES a naive call — every stem on `scripts/attention-rung-allowlist.txt`, which is to say the rows that allowlist exists to unblock. It runs in the required `agent-record` job, so #1545 alone would have turned `main` red. **The second was `assertGreater(excused, 0)`** in `test_the_ok_line_reports_the_excused_sites`. Same shape, one case down: `excused` counts unmarked calls in allowlisted files, so it reaches 0 when the LAST stem is cleaned up, redding the case while the checker is green at rc=0. It does not fire for any of the three rows individually, so it would have sat latent until the LTX-2.5 reroute tripped it. Both are repaired here rather than deferred, because none of this has landed: `scripts/check-attention-rung-consistency.py`, its allowlist and its suite are all CREATED by this pull request, so correcting a defective assertion in them is repairing the change, not amending a gate that `## Changing the rules or a checker` governs. The floor is now genuine non-emptiness (`>= 1`). One new case asserts every allowlisted stem NAMES AN EXISTING model source — keyed on file existence and deliberately never on scan membership, because a stem stops having a call site the moment its removing row lands, which is the state the allowlist is built to survive and which `stale_allowlist_entries` already promises in its own docstring ("Reported, never fatal"); asserting scan membership would have rebuilt the identical lock one line over. The `excused > 0` floor is gone, its shipped-tree half kept because it RE-DERIVES the count instead of pinning it and so holds at 0 as well as at 3, and the coverage it was standing in for moved onto two cases driven over a constructed scan and a temporary allowlist, which the model tree cannot switch off. RED first, each mutation proven applied and restored by sha256. Stubbing `scan_models` to `{}` reds the floor at `0 not greater than or equal to 1`, and independently so does renaming the checker's regex. A typo'd `muse_glimmer_vison` entry reds the new case naming that stem while the CHECKER stays green — which is the point, since a bogus stem is reported only as STALE and never fails. For the second lock the control is sharper: with the checker's excused count broken AND the tree in its end state, the retained shipped-tree case goes GREEN over the defective checker and only the constructed case catches it, so the replacement is real coverage rather than a deletion in disguise. Composed green, measured on this head with #1579's `muse_glimmer_vision.cpp` copied in: with the allowlist stem left in place the suite is `Ran 34 tests ... OK` and the checker prints `STALE (not a failure)` at rc=0, where before the repair it read `8 not greater than or equal to 9`. With the stem also deleted, only `test_allowlist_holds_only_the_in_flight_stems` reds, which is the by-design pin on the set, and updating that set in the same change returns it to green — so both removal routes now have one. A fresh review of the repair returned one finding, repaired here, and repairing it turned up a second of the same kind. The new case's comment claimed it also caught a checker printing `sites - marked`. It does not, and no constructed scan could make it: `main()` reaches the OK line only when `drift_sites` is empty, and then every unmarked site is excused, so the two quantities coincide identically — 1000 reachable green states enumerated, 0 where they differ. Substituting `sites - marked` into the checker leaves all 34 cases green. Nearby, `test_the_excused_count_is_not_sites_minus_marked` claimed that dropping the allowlisted file's marked site made the two diverge; it does not, because that lowers `sites` and `marked` together. Both comments now state what is actually pinned and why the rest cannot be, which is the same "a comment claimed more than the code delivers" class this branch already repaired twice. The fix is comment-only, and that is proven rather than asserted: `ast.dump` is byte-identical across the change, so no assertion, fixture or docstring moved. Five rounds of comment repair were needed, because each of the first three removed a false claim by writing a NEW causal explanation that the next fresh review then measured false. From the third round on the rule was to DELETE rather than re-explain, and to run every clause left standing. Comment lines go down, not up. `ast.dump` is byte-identical across rounds three and four, so no assertion or fixture moved; round five changes exactly one `assertTrue` MESSAGE string, with the assertion condition's own `ast.dump` hash shown identical either side. Two findings from those reviews were REFUTED by measurement rather than applied. `test_widening_the_regex_to_the_fast_rungs_is_visible` was reported dead because it survives a `drift_sites` break and a `\b` removal — but neither is the widening it names, and mutating `_NAIVE_CALL` to `\bvt::Attention\w*\s*\(` reds it along with five others. The spec's "six mutations that must go RED" was reported as five; there are six, and each reds under the mutation it names. Applying either would have renamed a working test and made an accurate record false. **Knowingly shipped, and recorded rather than hidden.** Three comments in `scripts/check-attention-rung-consistency.py` (`:58-61`, `:93-96`, `:252-255` — anchors measured, not estimated) are MEASURED FALSE and are NOT repaired here. All three are now enumerated in the test file beside the case that pins the real behaviour, so a reader of the suite can find every one. They claim the `\b` in `r"\bvt::Attention\s*\("` is what stops the pattern matching the fast rungs. It is not: the trailing `\(` does that, `vt::AttentionDenseFlash(` matches with neither, and the two patterns differ on exactly the 63 identifier characters — `xyvt::Attention(` alone. With the `\b` removed the suite stays green at `Ran 34 tests ... OK`, so the companion claim that the widening "is caught in this suite" is false too. The equivalent claims in the TEST file ARE repaired, because `tests/scripts/test_*.py` is not a governance checker. **The tree therefore contradicts itself between the test and the checker beside it, and that is deliberate**: `scripts/check-pr-size.py:170` classifies every `scripts/check-*.py` as a `governance_checker` and demands executable red-before evidence, which a comment-only diff cannot produce. Attempting the repair returns `ERROR: BASE checker stayed green ... changed test is not semantic evidence` at rc=1. The prepared patch was deliberately not committed rather than land a red `pr-size`. #1629 records both drift locks. #1631 records the comment freeze, and it is not one file: the pattern covers all 43 `scripts/check-*.py` checkers plus the `.sh` ones, so a false comment in any checker in this repository cannot be corrected on its own today. Both are linked in the three places AGENTS.md requires rather than in this body alone: `.agents/issue-index.md` gains one appended row each, and `## Owed` in `.agents/specs/attention-rung-visibility.md` records #1629 as DISCHARGED HERE and #1631 as owed with the reason it cannot be. No gate would have caught their absence, because `check-agent-record.py` counts index rows that name no owner and there was no row at all. ## Owed, and named rather than skipped Nothing on a CPU-only box executes `LaunchAttentionDenseFlash`, so the pure-arithmetic cases stay green over a launcher that lost the bound. The on-device refusal case emits a loud PENDING message and returns, and #1573 owns running it plus the reachability mutation that proves the case reaches the guard. No lease was taken: `dgx:gpu0` was unavailable for the whole branch. The 101,376-byte GB10 ceiling above does NOT discharge #1573 — it bounds what an opt-in could buy and says nothing about whether the launcher's refusal executes. Owed and filed rather than left to be discovered: **#1629**. The new `test_the_population_is_not_empty` case asserts the scanned `vt::Attention` population is `>= 9`, and the shipped tree has exactly 9 sites, so the floor has ZERO headroom and any row that deletes a naive call reds it. That is every stem on `scripts/attention-rung-allowlist.txt` — `muse_glimmer_vision` (#1545) and the two LTX-2.5 files — which is to say the rows the allowlist exists to unblock. Measured while landing this change, with #1579's `muse_glimmer_vision.cpp` copied onto this head: leaving the stem reds that case at `8 not greater than or equal to 9`, and deleting the stem reds it AND `test_allowlist_holds_only_the_in_flight_stems`, so a removing row has no green path. The tree was restored byte-for-byte after each and the suite returns to `31 tests ... OK`. This is the same drift-lock shape the kernel-matrix cell above was corrected for, retained one file away, and it is NOT repaired here because changing the floor changes what the gate accepts and AGENTS.md routes that to its own row, spec and red-before evidence. #1629 carries the evidence and two candidate directions, and #1579 is held on it. Inherited red, not introduced here: `test_cpu_x86_llamacpp_floor` fails in `agent-preflight.sh` on this box. It is the known load-dependent case of #618 — at high loadavg the harness exits `NO_QUIET_WINDOW (4)` where the case expects `GIVING_UP (2)` — and both the case and `scripts/cpu-x86-llamacpp-floor.sh` are byte-identical to `origin/main` on this branch, which is how it was established as inherited rather than assumed to be. FIVE CI jobs are red on the head, every one of them inherited from main with a named owner, and each was verified per-job against a main baseline rather than asserted. `windows-msvc-cpu` and `windows-msvc-vulkan` are the standing PR-only red (#584, #965). `build-test-cpu` and both `sanitize-cpu` arms fail on ONE shared doctest case, `test_runner.cpp:1557`; that case landed on main in `e2a9e035d` (#1273) and is owned by #1608 and #1602. Inheritance was established by comparing the failing ASSERTION and not the job name: the scheduled main baseline at `e2a9e035d` fails the identical `CHECK_THROWS_WITH_AS( make_runner(), "Block size must be a multiple of 16", std::invalid_argument )` at the same `test_runner.cpp:1557`, with the same "No valid attention backend for device type 0" text, in all three jobs, and with zero sanitizer findings in either sanitized arm. `test_runner.cpp` is not touched by this branch. `build-newest-gcc` was the sixth red when this body was first written and is GREEN here. #1581 and #1618 landed the `::getpid` repair on main, and this branch was re-merged onto `2e7f3bee7` to pick it up, so the job now compiles and reports on this change. Every job that can see this change is green, including `pr-size` and `agent-record`, which is the job that runs this row's new checker. Closes #1544. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
localai-bot
added a commit
that referenced
this pull request
Aug 22, 2026
… handing out an address-keyed cache (#1415) Row `A2-Q2b`, issue [#810](#810), spec `.agents/specs/nemotron-h-a2q2b-realckpt-lmhead.md`. Measured first, built second. The row's premise was that NemotronH decode re-expands quantized weights on the host and that `lm_head` is a large share of it. That was arithmetic, and the brief that dispatched this row said so. It is now a count. **This description was rewritten after a fresh review returned FINDINGS.** The measurement and the seam design below survived review unchanged and were verified independently. Three claims that were in the previous description did not, and they are corrected in place rather than quietly dropped: the new test had never compiled, the routing checker was said to hold a branch it does not hold, and the Marlin arm was said to be unbuildable here when it builds in seconds. **A scoped re-review then found no correctness defect in the shipped code and verified all seven repairs, and a final round corrected what it did find — all of it evidence, records and process.** The evidence tables named a tree they were not measured on and now carry an `overlay` column with the un-overlaid baseline stated as its own row; the Marlin object figure is re-measured and anchored to a SHA; the shared process-static this row inherits is named; the #1421 index row's counts are corrected while the pull request is still unmerged and the row still editable; and `§4` of the spec stopped telling the next agent to take `flock` on a fleet device. No product code changed in that round. ## What the measurement says Tallied at `NemotronHOwned::DenseBf16` (`nemotron_h.cpp:395`), the single seam every host dequant in this architecture passes through, on the real 21 GB checkpoint (revision `29f2d174`) through the production ABI driver `examples/nemotron_h_gen`. One decode step, T=1, `top_k` 6, 23 MoE layers, 369 dequant calls: | group | shape | calls | elements | % | |---|---|---|---|---| | routed expert `up_proj` | `[1856, 2688]` | 138 | 688 472 064 | 22.36% | | routed expert `down_proj` | `[2688, 1856]` | 138 | 688 472 064 | 22.36% | | shared expert `down_proj` | `[2688, 3712]` | 23 | 229 490 688 | 7.45% | | shared expert `up_proj` | `[3712, 2688]` | 23 | 229 490 688 | 7.45% | | `lm_head` | `[131072, 2688]` | 1 | 352 321 536 | 11.44% | | mamba `out_proj` (FP8) | `[2688, 4096]` | 23 | 253 231 104 | 8.23% | | mamba `in_proj` (FP8) | `[10304, 2688]` | 23 | 637 034 496 | 20.69% | | **TOTAL** | | **369** | **3 078 512 640** | **100%** | `138 == 6 x 23` exactly, which is what identifies this as the decode shape rather than a prefill aggregate. Three findings, and the first two are corrections: 1. The dispatching estimate put `lm_head` at 131072 x 4096 = 537e6 elements and ~43% of the population. `hidden_size` is 2688. It is 352 321 536 elements and 28.35%. 2. The widely quoted "1.24e9 elements / 2.49 GB per token" names a REGIME, not a total. It is not the host arm's 3 078 512 640. It is `mamba + lm_head` = 1 242 587 136 = 2.485 GB, the residue after A2-Q2a moved the MoE arm to the device, matching to four significant figures. 3. `lm_head` is the last one. Against the A2-D1 discriminator on `dgx:gpu0` (device mamba ON 1.554 s/token and 108.2x vs vLLM, OFF 10.319 s/token and 718.1x), the mamba arm is worth 6.64x and `lm_head` is on the host in every leg. Once the mamba arm lands, `lm_head` is 352 321 536 of 352 321 536, i.e. everything that remains. It is also the largest single re-expansion in the model by 12.7x, so its 704.6 MB transient bf16 buffer is the one that matters on a unified-memory box that reboots rather than OOM-kills. So the estimate's share was wrong, its direction was right, and the case is stronger after the discriminator than before it. ## The residency decision, taken explicitly The spec's `## 5. Owed` required this row to choose rather than default: "A2-Q2a routed around it by never calling either `MarlinDenseResidentFor`; `lm_head` must do the same or say why not." That accessor keys its repack cache on the weight's ADDRESS ([#984](#984)), and NemotronH is exactly the second-consumer condition an address key cannot survive. Rather than fork the seam as `qwen3_5.cpp` did, or inherit the defect as routing through it plainly would, the seam is EXTENDED: `MatmulNvfp4W4A16D` and `MatmulNvfp4MarlinD` accept a caller-owned `MarlinDenseResident*`, defaulting to the existing cache so every current caller is byte-for-byte unchanged. NemotronH owns its resident in a `ResidentSlot` on the weights, which is the property [#237](#237) added `ResidentSlot` for and the same call A2-Q2a made for the MoE arena. [#984](#984) is left exactly as it was for every other caller, where its own two-engine red-before belongs. ## What moved `NemotronHPagedForward`, the production path, now returns a device-resident `ForwardLogits` on an NVFP4 checkpoint, which is every released one, and downloads `final_normed` only when a consumer needs it. `NemotronHDeviceForward` deliberately keeps the host projection. It is the operand the numeric gate compares against, it is what makes A2-R's token gate attributable, and it has no production caller. The spec required this row to disclose where that property ends: it ends in the paged forward. ## The allowlist entry is narrowed, not removed Its one clause is met. But `check-runner-routing-consistency.py` resolves a cross-TU delegate only for the `Class::ForwardDevice` shape, and NemotronH's device forward is a free function in another TU, so the model still classifies HOST although `NemotronHPagedForward` assigns both `fl.device_tensor` and `fl.device_storage`. Filed as [#1410](#1410) and left to its own row: it changes checker semantics, which `AGENTS.md` routes away from an in-flow fix, and widening a classification to turn a red gate green is precisely the move that section slows down. ## The review repair ### The new test had never compiled, so the red it claimed could not have existed `tests/vllm/models/test_nemotron_h_moe_device.cpp` used `NemotronHHostWeights` unqualified and never declared it. Measured on the same command: | tree | `compile_rc` | errors | |---|---|---| | PR head `29b1128e`, this file | 1 | 9 | | merge base, same file | 0 | 0 | | repaired, default arm | 0 | 0 | | repaired, `-DVT_MARLIN_NVFP4=1` | 0 | 0 | The consequence is larger than the declaration. The file had never compiled on any platform, so **the red-first result this row claimed for its synthetic numeric gate did not exist and could not have existed**. That gate is CUDA-only and still has never executed; the spec now records `NEVER RUN` rather than restating a red nobody saw. The test target is deliberately registered with no CUDA guard, and that is what surfaced this — a case that skips at run time still has to parse on every CPU build. ### The production source asserted a protection that does not exist A comment at the device branch claimed the allowlist entry was removed and that "the routing checker, not a comment, is what now holds this branch in place: put the host projection back and it goes red by name". All three claims were false, and they contradicted this description and the allowlist file in the same commit. Reproduced here on the repaired tree: baseline `check-runner-routing-consistency.py` exits 0 with "3 host-logits off-framework (3 allowlisted)"; after deleting the whole `if (DeviceLmHeadEligible(...)) { ... return fl; }` block (`git diff` moves by 1 insertion and 14 deletions, so it applied) the checker exits 0 with **byte-identical output**; the file was restored to an identical sha256. **Nothing automated holds that branch.** The checker is not widened to make it — that changes checker semantics and [#1410](#1410) owns it with its own red-before. The comment now says what is true, the allowlist entry's own text is corrected the same way, and the reachability deletion mutation is recorded PENDING a CUDA window rather than claimed. Every case that can enter the device branch is `MarlinW4A16Selects`-gated and therefore CUDA-only, so that mutation cannot run on a CPU box. ### A silent dequant fallback that no token gate can see `DeviceLmHeadEligible` restated the shared dispatcher's three selection clauses and dropped `MarlinW4A16Enabled()`. Under an explicit `VT_NVFP4_MARLIN=0` the model said eligible while `MatmulNvfp4W4A16D` took its naive redundant-dequant arm — the same logits, while re-uploading the whole `[131072, 2688]` operand on **every decode step**, because `LmHeadNvfp4View` hands out a stack temporary that `ResidentNvfp4`'s weight-keyed cache can never hit. The repair is structural rather than a patched clause. The three clauses now live once, in `dense_nvfp4::MarlinW4A16Selects`, and both the dispatcher and the model call it, so they cannot drift. `DeviceLmHeadD` additionally refuses BY NAME if that same predicate is false on the operand it is about to hand over, which catches an eligibility answer taken against a different queue or dtype. That is deliberately the predicate and not the seam's `fallback_gemms` counter: `MutableW4A16Stats()` is a plain non-atomic process-wide static, so a counter window in production would refuse a correct run whenever anything else took a fallback GEMM concurrently, and a false refusal is worse than the silence it replaces. The counter is the right instrument in a test, where single-threadedness is a property of the harness, and the synthetic gate asserts it there. It is demonstrably armed rather than assumed: `test_qwen3_forward.cpp:559` — the assertion itself, not the `TEST_CASE(` line at `:497` — already asserts on CPU that it reaches exactly `5 * num_hidden_layers` when the dispatcher does fall back. **And the counter is not the only process-wide static on this route, so this description should say what the row does NOT decline.** `dense_nvfp4::DenseMarlinWorkspace` (`dense_nvfp4_gemm.h:506`) is a process-static DEVICE allocation — `static void* ws` behind a `static std::mutex`, sized from `MarlinDeviceSms` and keyed on nothing, not even the device index — shared by every caller of `MatmulNvfp4MarlinD` (`:539`) and `GateUpFusedMarlinD` (`:700`). It has three consumers today (the shared dense route in `dense_attn_block.h`, MiniMax-H3, and the compressed-tensors NVFP4 scheme); NemotronH's `lm_head` becomes a fourth. It is PRE-EXISTING — it arrived at `80d1da096` and its definition is byte-identical at this row's merge base and at its head — and nothing here addresses it. In a change whose thesis is refusing to inherit the sibling process-static defect ([#984](#984)), saying nothing about the one it does inherit was an asymmetry a reviewer is entitled to see stated. The cost the fallback would carry is 198.18 MB per call — 176.16 MB of packed codes plus 22.02 MB of group scales — arithmetic that checks to the byte against the `[131072, 2688]` geometry. The behavioural red for this class needs CUDA and is PENDING with the rest. ### The Marlin arm builds here, and saying otherwise overstated the blocker The previous description said "this branch has not built the Marlin path at all". `include/vt/cuda/marlin_repack.h` includes only `<cstdint>`, `<cstddef>` and `<vector>`, so the host side needs no CUDA toolkit. On this box, which has no `nvcc`: ``` c++ -std=c++20 -I include -I src -isystem third_party -DVT_MARLIN_NVFP4=1 \ -Wall -Wextra -Werror -c -o nhd_marlin.o \ src/vllm/model_executor/models/nemotron_h_device.cpp -> rc 0, 0 errors, 0 warnings, a 1 268 552-byte object ``` **That byte count is anchored to the merge commit on this branch**, gcc 13.3.0, no `nvcc`, and compilation is deterministic here — two runs of the identical command, identical `sha256 42d670b6...`. The anchor is the correction. The figure this description carried until now, 1 269 696, was accurate when written and accurate at the reviewed head `29b1128e3`, but the translation unit moved twice afterwards — 1 272 808 at `fedf78d86`, 1 268 552 from `bff2b7b2f` on — and an evidence block naming no SHA cannot tell a reader which of the three it means. The measured series across seven commits is in the spec. The size is incidental either way; what the block asserts is `rc 0`, zero errors and zero warnings on a toolchain with no CUDA, and that holds at every one of those commits. Both arms are compiled and both are compiled `-Werror`. What genuinely needs `nvcc` is the Marlin **kernel** and every **execution** of the device path; those stay PENDING and are what the table below records. ### Reachability, and a CPU vehicle that already existed Both synthetic cases build `NemotronHHostWeights` by hand, which proves the class works and not that anything reaches it, and the deletion mutation for the device arm needs CUDA. But this row also refactored the part of `NemotronHPagedForward` that every CPU step runs, and `tests/vllm/models/test_nemotron_h_paged_forward.cpp` already drives that function through a real `GPUModelRunner`. It gains a case: three gathered rows against one request, so the returned row count must follow `logits_indices` and never `num_reqs`; the trace's copy of `final_normed` fed back through the production host entry point must reproduce the returned logits bit for bit; and the same logits must come back with the trace off. With `trace->capture` set, `DownloadF32(final_normed)` ran twice on the host arm. One download now serves both consumers on every path. That duplicate is not observable from outside the function — both copies are of the same unchanged buffer — so it is repaired structurally and no assertion pretends to catch it. ### Records `spec §4.3` did not exist; the requirement is under `## 5. Owed`. The allowlist's `## Gates` did not either; the heading is `## 3. The gate`. `docs/FEATURES.md` still listed the device `lm_head` as owed while `docs/USAGE.md` said it was device; both now say implemented and never executed, which is the honest state. The appended issue-index row for [#1410](#1410) cited `## Owed` where the heading is `## 5. Owed` — cosmetic, since the row passes on its owning-row disjunct, but the index is append-only and this is the only window to correct it. [#1421](#1421) is filed rather than fixed in flow: two NVFP4 W4A16 upstream anchors are wrong at the pin, verified against a checkout of `vllm-project/vllm` at `5559679229bc961848b121ccdeaa8fa5d79bec98`. `kernels/linear/__init__.py:879-881` is the tail of `init_wfp8_a16_linear_kernel`, an FP8 helper sharing the `a16` vocabulary, and the forced-Marlin line it quotes is at 922-924; `qwen3.py:271-274` is the `class Qwen3ForCausalLM(...)` line and `packed_modules_mapping` is at 275-278. Both predate the merge base, and both are cited from nine code and test sites over SEVEN files plus eight record sites over FIVE record files (`parity-ledger.md` carries three). The index row said "six files" while its own parenthetical enumerated seven, and "six record files" where there are five; `.agents/issue-index.md` is append-only and a landed row can never be edited, so an unmerged pull request is the only window to correct that, and this one does. The set still needs re-deriving at the pin rather than a sed. This branch relocates one of those citations and deliberately leaves its VALUE unchanged, so the issue owns every site uniformly instead of half of them looking reviewed. ## Gate status, stated as pending rather than as passed Every number here is from a run on a CPU box with no CUDA toolkit, `RelWithDebInfo`, `-Wall -Wextra -Werror`. **Read the `overlay` column first.** This tree cannot construct a `GPUModelRunner` at all — [#1371](#1371) throws there — so every green below was taken with [#1392](#1392 production fix applied to the working tree, never committed here and reverted byte-for-byte afterwards. The previous version of this table said "a run on this tree" and reported the §12 case at `run_rc 0 / SUCCESS!`, which is false of the tree it named: un-overlaid, that case is already red. | binary | overlay | `run_rc` | cases | assertions | verdict | |---|---|---|---|---|---| | `test_nemotron_h_moe_device` | none | 0 | 4, 4 passed | 4 | `SUCCESS!` — and **all four SKIP**, both A2-Q2b cases included | | `test_nemotron_h_paged_forward`, whole binary | **none** | 1 | 13, 2 passed, 11 failed | 18 | `FAILURE!` — [#1371](#1371), not this row | | `test_nemotron_h_paged_forward`, whole binary | **#1392** | 0 | 13, **13 passed** | 3269 | `SUCCESS!` | | the new §12 case alone | **none** | **1** | 1, **0 passed, 1 failed**, 12 skipped | **0** | **`FAILURE!` — throws #1371 before its first assertion** | | the new §12 case alone | **#1392** | 0 | 1, 1 passed, 12 skipped | 13 | `SUCCESS!` | **The `moe_device` row is not a pass, and it is the honest state of the synthetic gate.** The binary builds and exits 0, but every case takes the `TryCudaQueue` skip on a GPU-less box, so those 4 assertions are the skip notices themselves. The numeric gate examined nothing. **The paged-forward red is [#1371](#1371 All 11 failures throw the identical `No valid attention backend for device type 0 from {FLASH_ATTN: [head_size not supported]}` at `GPUModelRunner` construction, and 10 of them are cases this row never touched. Overlaying [#1392](#1392 production fix in the working tree — never committed here, reverted afterwards — turns the same binary green. Worth recording: doctest printed `assertions: 18 | 18 passed | 0 failed` while 11 cases were throwing, so the assertion line alone would have read as a pass. ### The red-first, on the cases that can run Each mutation went into a scratch copy of `nemotron_h_device.cpp`, was proven applied by `git diff --stat`, was built before being run (a mutation that fails to build proves nothing), and was restored to an identical sha256. **Every row below was also run with [#1392](#1392) overlaid**, and that is not a detail. Un-overlaid, the §12 case is already `run_rc=1` on the UNMUTATED tree, so an M1 or M2 red measured there would prove nothing whatever. The overlay is what lets the vehicle report a green in the first place; only then can a mutation take it away. | mutation | overlay | applied | `compile_rc` | `run_rc` | verdict | |---|---|---|---|---|---| | — (unmutated control) | **#1392** | — | 0 | 0 | **GREEN**, `SUCCESS!`, 13 assertions | | — (unmutated control) | **none** | — | 0 | **1** | already **RED** — #1371, which is why the rest is overlaid | | **M1** — `n_out` -> `R` at the host projection | **#1392** | 1 ins / 1 del | 0 | **1** | **RED**, `FAILURE!` | | **M2** — never fill `trace->final_normed` | **#1392** | 1 ins / 2 del | 0 | **1** | **RED**, `FAILURE!` | | **M4** — restore the exact pre-repair two-download shape | **#1392** | 3 ins / 8 del | 0 | 0 | **GREEN — reported, not hidden** | Re-derived rather than restated, on a clean tree at the merge commit, overlay and mutations reverted to byte-identical sources (`nemotron_h_device.cpp` back to `sha256 abf6e21f...`): §12 alone un-overlaid gives `run_rc=1`, `1 | 0 passed | 1 failed | 12 skipped`, `assertions: 0`, `THREW exception: No valid attention backend for device type 0 from {FLASH_ATTN: [head_size not supported]}`; overlaid, `run_rc=0`, 1 passed, 13 assertions; the whole binary overlaid, 13/13, 3269 assertions; M1 `compile_rc=0 run_rc=1` throwing `gathered row count does not match hidden_size` at `nemotron_h.cpp:1029`; M2 `compile_rc=0 run_rc=1` on `REQUIRE( 0 == 288 )`. M4's cell needs no separate attestation — un-overlaid every run of this case is red, so a GREEN is only reachable with the overlay and its verdict entails the column. M1 is the red the `n_out` rename exists for: with the request count substituted, the returned row count is 1 where the gather asked for 3. M2 arms the trace-operand assertion. Both showed the same trap — doctest printed `assertions: 2 | 2 passed | 0 failed` on M1 while the case was failing, because a `REQUIRE` throws rather than counting. M4 staying green is a result, not a gap. The duplicate `DownloadF32` copies the same unchanged buffer twice and produces identical bytes, so nothing observable from outside the function can distinguish it. It is repaired structurally and no assertion pretends to catch it. **The device arm's own red-first does not exist on a CPU box and is not claimed.** `MarlinW4A16Selects` is false on a CPU queue, so the device branch, the `fallback_gemms` assertion and the reachability deletion mutation are all unreachable here. ### A gate this row was already failing, and cannot repair in place `scripts/check-doc-checkpoint.py` is red on this branch, and it was red at the reviewed head `29b1128e` with the identical two errors (`--base 96ed834 --head 29b1128`): ``` ERROR: commit 1c62d99: changed user_usage but did not update docs/USAGE.md ERROR: commit 8fa900a: changed .agents/benchmark-record.md: measurement recorded but did not update docs/STATUS.md ``` The previous description said this gate had one real failure and that it was repaired. It was not, and `.github/workflows/ci.yml:519` runs the same `--base/--head` invocation, so the lane is red for this reason independently of [#1371](#1371). The second error names a real gap and it is now closed: this row moves a lifecycle state and records a measurement, and `docs/STATUS.md` said neither. It does now, in 192 characters, because that page carries a shrink-only ratchet on oversized cells and the first attempt took it from 44 to 45. The first error, and the historical form of the second, cannot be closed by a later commit. The checker iterates `commits_in_range` and judges each commit on its own contents, so the obligation belongs to `1c62d9974` and `8fa900a62`. Discharging it means rewriting commits that are already the reviewed base, which resets this pull request's CI approval and moves the head a fresh reviewer was asked to look at. **That is a scheduling decision, so it is recorded as owed and raised for the operator rather than taken by a repair pass.** Run over this repair's own commits (`origin/main..HEAD`), the checker reports these two and nothing else. ### The spec sent the next GPU window through the wrong mutex `§4` step 1 read "Take `$GPU_LOCK` with a blocking `flock` and wait". `AGENTS.md` requires a fleet device to be claimed through `rc`, and `dgx:gpu0` is one. The text predates that rule, but this branch is the wrong place to leave it standing: it ADDS `scripts/nemotron-h-a2q2b-gpu-gate.sh` for exactly that window, the script correctly takes no mutex of its own, and the contradicting instruction sat beside it in the same file — and every PENDING leg above is waiting on that window, so the next reader of `§4` is the person about to open it. Two mutexes that do not exclude each other are worse than one: the fleet cannot see `$GPU_LOCK`, so a `flock` over `ssh` does not exclude a concurrent `rc` holder and the controller reports the box free while somebody is on it. On 2026-08-17 that pair voided a whole speed axis (`minimax-music3.md` §13.10). Step 1 now claims through `rc run -d dgx:gpu0 --max-runtime <N>h --` with the gate script as the payload, and step 2's headroom check moves inside the lease and names the script's own PRECONDITION 1. ### The trailer gates had stopped running At the previous head, `origin/main` was not an ancestor, so `scripts/agent-preflight.sh --fail-on-skip` took its `TRAILER_BEHIND` arm and SKIPPED both `commit-trailers` and `commit-style`. A skip is not a green. **It then happened five more times during the device-leakage repair**, because `origin/main` moved under it at `9ecaf1bb3`, `f07f96e1c`, `aeba0de6f`, `01854663c` and `c8d926ea8`. The ref only moves on a fetch, so the fix is ordering: merge, then do not fetch again, then gate, then push. At the pushed head both gates RUN: `BASE_SHA=c8d926ea8` non-empty, `ANCESTRY_STATUS=0`, `RANGE_COUNT=22`, `RANGE_STATUS=0`, **zero `SKIP` tokens anywhere in the report**, 85 gates `ok`, and the only failure is `doc-checkpoint range` with the two commits above. The skip arms are falsified individually rather than read off a green line, because `agent-preflight.sh` prints a summary only on failure. Each is armed over exactly this range, proven by a detached scratch commit with a period-terminated subject and no trailer block, which turns `check-commit-style.py` and `check-commit-trailers.py` red with the specific messages. Both files that both sides of the merge touched are records, resolved by shape rather than by whatever the three-way merge produced. `.agents/issue-index.md` is a genuine append-only log carrying `merge=union`, and both sides append at the tail, so all four new rows survive — verified additive, 469 + 2 lines on each side to 471, no line removed or altered. `docs/USAGE.md` is a keyed record: the merged file is byte-identical to `origin/main`'s version with this branch's scoped edit re-applied, so every key neither side owns is unchanged. ### The DSR ratchet was red, and no review round ever saw it run `device-leakage` never COMPLETED while this pull request was under review, so its verdict was an input to none of the three passes. It completed after the third and failed: `vt_ifdef` **35 against a baseline of 32**, `rc 1`. Three `#ifdef VT_MARLIN_NVFP4` sites had been added to the device-agnostic shared layer, which is exactly what the ratchet exists to stop. They were located by running `scripts/check-device-leakage.py --report` at the failing head `7a3909187` and diffing the per-file table against `9ecaf1bb3`, rather than by reading the diff for guards. | # | site at `7a3909187` | what the guard decided | resolution | |---|---|---|---| | 1 | `include/vllm/model_executor/models/dense_nvfp4_gemm.h:768` — inside `MarlinW4A16Selects` | **nothing** | **removed** | | 2 | `src/vllm/model_executor/models/nemotron_h_device.cpp:883` — around `LmHeadNvfp4View` | **nothing** | **removed** | | 3 | `src/vllm/model_executor/models/nemotron_h_device.cpp:985` — `DeviceLmHeadD`'s body | `ResidentIn` and a complete `dense_nvfp4::MarlinDenseResident` | **`DSR-ALLOW(A2-Q2b)`** | **No baseline was changed.** `AGENTS.md` forbids making a red gate green by widening an assertion, and a baseline bump is that. `scripts/device-leakage-baseline.json` is untouched at 32, and the per-file table is now byte-identical to `origin/main`'s. **(1) is the case the checker's own message describes.** `MarlinW4A16Selects` is a SELECTION wearing a build guard, and every term it reads exists in every build: `MarlinW4A16Enabled()` is declared above the guarded region, and `vt::OpRegistered` is the op/provider table's own answer to whether the Marlin arm is realized for a device. The flag and the registration are one condition, not two — `CMakeLists.txt`'s single `if(VLLM_CPP_MARLIN)` block adds `src/vt/cuda/cuda_moe_marlin.cu`, whose file-scope `Registrar` holds the tree's only `RegisterOp(OpId::kMoeGroupedGemmNvfp4Marlin, …)`, and defines `VT_MARLIN_NVFP4=1` in that same block. A build without the macro registers nothing, so the query already resolves false on exactly the builds the `#ifdef` excluded. This is the call `nemotron_h_device.cpp`'s `moe_on_device` selection had already made, in a comment that says so. **(2) was measured rather than reasoned.** `LmHeadNvfp4View` names nothing the Marlin build adds — `Nvfp4Weight` comes from `qwen3_5_weights.h` and `OwnedBytes` from the loader, both unconditional, and no `vt::cuda::` symbol appears in it. The claim that its external linkage at `namespace vllm` scope is what makes an unused definition harmless where its only call site is compiled out is proven by the mutation that removes that property: adding `static` turns the same CPU compile RED at `rc 1`, `error: 'vllm::Nvfp4Weight vllm::LmHeadNvfp4View(...)' defined but not used [-Werror=unused-function]`. The file was restored to an identical sha256 (`9719ea70…`) afterwards. **(3) is TYPES-not-behaviour and takes the checker's documented escape hatch.** `DSR-ALLOW` is not a baseline change: the site is excluded from the count but COUNTED AND PRINTED on every run, so the exemption is visible in CI output rather than invisible in the diff. `DeviceLmHeadD`'s body names two symbols that do not EXIST without the guarded arena region — the `ResidentIn` template, defined inside it, and `dense_nvfp4::MarlinDenseResident`, which the header declares unconditionally and defines only under `VT_MARLIN_NVFP4`, so the reference cannot bind to an incomplete type. It is the same class and the same stated reason as the five sibling guards A2-Q2a and A2-P already carry in this file, and its `#else` refuses by name. The SELECTION for this arm stays a runtime op-table query. **Measured on this tree**, a CPU build with `VT_MARLIN_NVFP4` absent from `build/compile_commands.json` — positive control: 1020 `VLLM_CPP` hits in the same file, so the grep is not silently wrong — which is the configuration that exercises both removals, because it is the arm the deleted `#else` branches used to serve. | what | before (`7a3909187`) | after | `rc` | |---|---|---|---| | `check-device-leakage.py` `vt_ifdef` | 35 | **32** | 1 → **0** | | `DSR-ALLOW` exemptions in force | 20 | **21** | — | | `scripts/device-leakage-baseline.json` | 32 | **32, untouched** | — | | per-file table vs `origin/main` | +1 header, +2 model TU | **identical** | — | | `nemotron_h_device.cpp`, `nemotron_h.cpp`, `qwen3_5.cpp` at `-Wall -Wextra -Werror` | — | compile | **0** | Each repair is individually load-bearing, proven by reverting it alone in a scratch worktree and re-running the gate. Every mutation was verified applied by `git diff --stat` and restored to an identical sha256, with the unmutated control green immediately before and after. | mutation | applied | `vt_ifdef` | `rc` | verdict | |---|---|---|---|---| | — (control) | — | 32 | 0 | `ratchet holds` | | **M-B** — restore the guard on `MarlinW4A16Selects` | 4 ins | **33** | **1** | **RED**, `DSR REGRESSION` | | **M-C** — restore the guard around `LmHeadNvfp4View` | 2 ins | **33** | **1** | **RED**, `DSR REGRESSION` | | **M-D** — delete the `DSR-ALLOW(A2-Q2b)` line | 1 del | **33** | **1** | **RED**, `DSR REGRESSION` | | — (control, after restore) | — | 32 | 0 | `ratchet holds` | The two pre-existing allowlist entries the report also prints, `deepseek_v4_device.cpp [kcuda] x8` and `platform.cpp [dev_cast] x1`, are byte-identical at `9ecaf1bb3` and here. This change moves one bucket and nothing else. **`docs/USAGE.md` rides in the same commit, because it has to and because it is true.** Any commit touching `include/vllm/` is `user_usage` to `scripts/check-doc-checkpoint.py` and owes the surface in that same commit. `doc-checkpoint range` is already red on this branch for `1c62d9974` and `8fa900a62`, which an operator decision covers; a third error of the identical kind would be new damage rather than inherited, so the header hunk was recommitted with the USAGE edit beside it. The edit is not written to feed the gate: the NemotronH arms table enumerates what selects the host `lm_head` projection and omitted `VT_NVFP4_MARLIN=0`, a user-settable knob that selects an arm and appears nowhere else in USAGE, and this change is precisely what makes that knob reach the model's eligibility test in every build rather than only where the guard compiled it in. `check-doc-checkpoint.py --commit` is `OK` on that commit, with the still-red `1c62d9974` as the positive control that the checker is armed. **`origin/main` moved four times during this repair** — `9ecaf1bb3`, `f07f96e1c`, `aeba0de6f`, `01854663c` — and each time `git merge-base --is-ancestor` answered 1, which is the arm that SKIPS both trailer gates without failing. Each was merged rather than worked around, and every auto-merged keyed record was resolved by key on every merge: the branch delta across the merge byte-identical to the branch delta from the merge base, checked file by file with a positive control on the comparison itself. ### The three red CPU jobs were inherited, and the merge cleared them `build-test-cpu` and both `sanitize-cpu` arms failed on this pull request, all three on `test_nemotron_h_paged_forward` — [#1371](#1371), which [#1392](#1392) fixed on `main` after the review rounds. `sanitize-cpu` is NOT a no-baseline job, so that had to be triaged by WHICH test failed rather than assumed. The prediction is verified rather than asserted. With `origin/main` merged, the same binary is green on this tree with NO overlay: `13 | 13 passed | 0 failed`, `assertions: 3269 | 3269 passed | 0 failed`, `Status: SUCCESS!`, `run_rc 0` — the same numbers the spec's tables previously recorded only WITH #1392 applied to the working tree and reverted. **The `#1392` overlay is therefore retired**; the historical rows in the spec keep their `overlay` cells, because they describe the tree they were measured on and rewriting them would make them false. ### `main` brought four red suites with it, and they are attributed both ways The third sync carried `4712dac40` (`VT-ACT-ROUND-POLARITY`, [#1322](#1322) via [#1347](#1347)), and with it four red CPU suites this branch does not touch. Both `ctest` runs are on this box, no overlay of any kind. | tree | base | `ctest` | result | |---|---|---|---| | this row + `9ecaf1bb3` | before `4712dac40` | 567 | **`100% tests passed, 0 tests failed out of 567`**, `rc 0` | | this row + `01854663c` | after `4712dac40` | 569 | `rc 8`, **4 failed**: `test_minimax_music3_ar`, `test_ltx2_text_encoder`, `test_muse_glimmer_text`, `test_muse_glimmer_text_fallback` | Those four are [#1458](#1458), filed from another flow before this control was run. Inherited, not caused, and proven in both directions on the same tree rather than argued — each mutation verified applied by `git diff --stat` and restored to an identical sha256: | control | change | `compile_rc` | `run_rc` (music3 / ltx2 / glimmer) | verdict | |---|---|---|---|---| | **A** | revert THIS row's two source files to their pre-repair `7a3909187` content | 0 | **1 / 1 / 1** | still RED — **not this row's** | | **B** | revert `src/vt/cpu/cpu_ops.cpp` alone to `4712dac40^` | 0 | **0 / 0 / 0**, `37/37`, `27/27`, `24/24` | GREEN — **`4712dac40` is the cause** | Control A answers the attribution question on its own; control B is here because naming a cause is more useful to the next reader than clearing oneself. Neither repairs anything, deliberately: #1458 needs `VT-ACT-ROUND-POLARITY` to decide whether its kernel or four bf16 error floors that were never re-derived are the wrong side, which is that row's oracle work rather than a small and clear in-flow fix. **That window has since closed.** [#1458](#1458) was fixed on `main` and closed, and this branch has merged the repair, so the four are not expected on the current head. The controls stay recorded because a red that was attributed and then disappeared is still the reason two `ctest` runs in this description disagree, and deleting the evidence would leave that unexplained. ### `documentation-checkpoint` is the same two commits, and it is not new It went red for the same reason `device-leakage` did: it had never COMPLETED during review. Its whole failure set is `1c62d9974` and `8fa900a62`, the two an operator decision covers and which cannot be repaired in place — `check-doc-checkpoint.py` iterates `commits_in_range` and judges each commit on its own contents, so a later commit cannot close either, and discharging them means rewriting the reviewed base. The job's three invocations were reproduced locally at this head with CI's own arguments (`--base` the PR base `63d87805c`, `--head` HEAD): | invocation | result | |---|---| | `check-doc-checkpoint.py` | `rc 1` — **exactly `1c62d9974` and `8fa900a62`, nothing else** | | `check-now-current.py` | `rc 0`, `OK: .agents/NOW.md is a current, in-budget resume digest` | | `check-role-discipline.py` | `rc 0`, `OK: every change on main arrived on a task branch` | **This repair's own commits are clean, and that is not an accident.** The first draft of the header fix added a THIRD error of the identical kind, and it is the reason `docs/USAGE.md` rides in the same commit as the `include/vllm/` hunk rather than in a follow-up: the checker's obligation is per-commit, so a follow-up cannot discharge it. `check-doc-checkpoint.py --commit 65ab066` is `OK`, with the still-red `1c62d9974` as the positive control that the checker is armed against exactly that class. ### Nothing ran at all for a whole push, and the cause was a conflict rather than a queue The push before this one produced **zero check runs** on `2c63cc87a` while other pull requests were starting normally. `gh pr view` named it: `mergeable=CONFLICTING`, `mergeStateStatus=DIRTY`. GitHub could not build `refs/pull/1415/merge`, so no `pull_request` workflow had anything to run against. A healthy Actions queue plus a branch with zero runs is the shape a conflict makes, and reading it as a slow queue would have cost the whole verdict. The conflict was `docs/STATUS.md` and `docs/USAGE.md`, and neither was resolved by picking a side, because both had been **relocated** on `main`. `USAGE.md` is now a 231-line hub whose per-model content moved under `docs/models/`, and `STATUS.md` is a 27-row surface summary with no per-model rows at all. So both were taken WHOLE from `main` and this row's scoped edits re-applied BY KEY at their new home, `docs/models/nemotron-3-5-lightning.md` — 22 insertions and exactly ONE deletion there, every unrelated key byte-identical. One of this row's edits is **dropped rather than re-applied**, deliberately, and the other was never dropped at all — the claim that both were is corrected below. The `STATUS.md` Nemotron row is genuinely gone: merged `main`'s `STATUS.md` is a 98-line surface summary that enumerates no model at all. `Nemotron` returns 0 hits there, and so do the positive controls `Kimi` and `Laguna`, while the file plainly greps. Re-adding it would re-create a row a relocation removed, which is the duplicate a second relocation makes. The `USAGE.md` `NemotronHForCausalLM` refusal row is the opposite case, and the earlier claim that it "no longer exists anywhere under `docs/`" was **wrong**: `aee6c48d6` (#1491) did not delete that row, it **RELOCATED** it to `docs/reference/model-loading.md` — a file that did not exist at this branch's pre-merge base, which is precisely why the earlier pass read it as gone. It is still on `main`, and this row now updates it. `docs/FEATURES.md` auto-merged and kept this row's cell — checked, not assumed. `check-conflict-markers.py` (#1450, new on `main`) reports 0 findings over 3820 tracked text files. Run locally from the ORIGINAL merge base, `check-doc-checkpoint.py` also reports `d995c52f0`, `af25bd251` and `995ed1ccd` — each verified an ancestor of `origin/main`, none in this branch's own range, all arrived with the merge. CI does not see those, because it resolves `--base` to the CURRENT merge base: **its whole failure set is `1c62d9974` and `8fa900a62`**, the two an operator decision covers, and nothing else. This row's own commits stay clean either way: `check-doc-checkpoint.py --commit` is `OK` on `65ab06636` and on the merge itself. ### Two public rows still said the head runs on the host, and this change is what made them false Following the relocation above to its consequence: `docs/reference/model-loading.md:188` still read "`lm_head` and FP8 Mamba2 projections run on the host". That sentence was true before this row and is false after it, which is the test for whose obligation it is. Left alone, the tree would land with `docs/models/nemotron-3-5-lightning.md` calling the head a device arm and this row calling it a host arm, on the same merge commit. Only the clause this row falsified moves. `lm_head` reaches the device **on the paged forward** and is carried as implemented-and-unmeasured, matching what the model recipe and the spec's `## Now` already say. The FP8 Mamba2 projections are a **different unit** (A2-Q1, #1289, still held DRAFT) and still run on the host, so that clause is kept and only re-worded to stop sharing a verb with the head. The pending-token-gate and no-GGUF clauses are untouched, because both are still accurate. Sweeping by CLAIM rather than by file, at the MERGED tree, found a **second** site the scoped review did not reach, because it is not a refusal row. The checkpoint registry at `docs/USAGE.md:230` lists per-checkpoint arms, and its "Supported arms" cell read "host FP8 Mamba2 and NVFP4 head". That "NVFP4 head" is this same `lm_head`, so the same test applies and the same one clause moves. That cell is the surface the "say which weights, and from where" rule owns, so a supported-arms cell putting the head on the host is exactly the contradiction this finding is about. The sweep carried its own controls, because a null grep only proves the terms wrong. Every `docs/**.md` and `README.md` line asserting a host placement for a head or for logits is 11 hits, 2 of them the sites repaired here (positive control); the same pipeline with a nonsense term returns 0 (negative control). The other 9 were read rather than pattern-matched, and none is stale: `BENCHMARKS.md:16` is this row's own host-re-expansion ATTRIBUTION and is what motivates the device arm; `BENCHMARKS.md:276` is Laguna GB10 ATS weight residency; `ENVIRONMENT.md:144,187,212` are Qwen3.5/Qwen3.6 levers; `server.md:88` is the CPU-only `prompt_logprobs` full-logits route on the shared runner, not this model; `model-loading.md:144` is the FP8 scalar-scale guard; and two are historical `docs/superpowers/plans/` documents, which are not projections. **No gate can catch either repair.** `check-doc-checkpoint.py` asks whether SOME `docs/USAGE.md` edit accompanied a commit of the changed class. It never asks whether a sentence inside it is still true. ### Every check has COMPLETED on the current head, and the remaining reds are all inherited `ci` run `32537115264` on `99f9f672a`: **25 checks, 0 pending** — 14 pass, 6 fail, 5 skipped by design on a pull request (`attest`, `baseline-summary`, `manifest`, `promote`, `publish`). | check | on `b7d89b43e` | on `99f9f672a` | whose | |---|---|---|---| | **`device-leakage`** | SUCCESS | **SUCCESS** | **this row's.** The restructure survives the 38-commit merge at `vt_ifdef` 32, and `scripts/device-leakage-baseline.json` is byte-identical to `main`, so the row adds no new allowance | | `build-newest-gcc` | FAILURE | **SUCCESS** | was **`main`'s**, and it cleared exactly as predicted once `main` carrying #1581 was merged. CI builds the MERGE commit, so the red came from the base side; #1565 is closed | | `build-test-cpu`, `sanitize-cpu (address,undefined)`, `sanitize-cpu (thread)` | SUCCESS | **FAILURE** | **`main`'s**, [#1608](#1608) and [#1602](#1602) — see below | | `build-test-vulkan`, `build-test-cpu-arm64`, `cuda-fat-build`, `cuda-arch-features`, `vulkan-spirv-freshness`, `agent-record`, `pr-size`, `commit-protocol-tag`, `last-gated-commit`, `plan`, `verify` x2 | SUCCESS | SUCCESS | — | | `documentation-checkpoint` | FAILURE | FAILURE | `1c62d9974` + `8fa900a62` ONLY — the two an operator decision covers, verified against the job log rather than assumed | | `windows-msvc-cpu`, `windows-msvc-vulkan` | FAILURE | FAILURE | the two no-baseline PR-only jobs, [#584](#584) | The three CPU legs all fail on **`282 - test_runner`**, and the failure is inherited rather than this row's. The failing case, `runner: initialize_kv_cache refuses a non-multiple-of-16 block size`, has **0** occurrences at `b7d89b43e` — the head where all three legs were SUCCESS — and **1** at the merge base `5453e571d`. It was introduced by `e2a9e035d` (#1273), which **is** among the 38 commits the merge brought in, and this row's range touches **0** files in that test (control: it does touch the two `tests/vllm/models/test_nemotron_h_*.cpp` files, so the query discriminates). It is also **not a data race**, which matters because `sanitize-cpu (thread)` is the lane that would expose one. All three legs report the SAME deterministic assertion — `test_runner.cpp:1557`, `20 test cases / 19 passed / 1 failed`, `544 assertions / 543 passed / 1 failed` — and **0** sanitizer findings in any of them (`WARNING: ThreadSanitizer` = 0, `ERROR: AddressSanitizer` / `runtime error:` = 0). A race cannot make a non-sanitized `build-test-cpu` fail identically. The assertion expects `"Block size must be a multiple of 16"` but the attention registry refuses first, with `"No valid attention backend for device type 0 from {CPU_ATTN: [block_size not supported], FLASH_ATTN: [block_size not supported]}"` — which is precisely what #1602 describes. Nothing is suppressed here and nothing is pragma'd, because there is nothing of this row's to suppress. ### Gate table | leg | state | |---|---| | seam extension, device arm, production wiring | DONE; built on BOTH arms, `-Werror`, 0 warnings | | the CPU-reachable half, gated through a real `GPUModelRunner` | DONE and RUN, with M1/M2 red-first | | routing allowlist narrowed, [#1410](#1410) filed | DONE | | synthetic device `lm_head` numeric gate | COMPILES; CUDA-only; **NEVER RUN** | | `nvcc` build of the Marlin kernel, real-checkpoint numeric leg, token identity, reachability deletion mutation, the `VT_NVFP4_MARLIN=0` fallback red | PENDING a `dgx:gpu0` window, job queued | A reviewer should treat every CUDA leg as unrun. The spec's `## 6. Now` carries the same tables, so the pending state lives with the row rather than only in this description. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
localai-bot
added a commit
that referenced
this pull request
Aug 22, 2026
… the correctness-grade attention kernel (#1545) (#1579) `src/vllm/model_executor/models/muse_glimmer_vision.cpp:639` called `vt::Attention` for every segment of every one of the perception encoder's 50 blocks. `src/vt/ops.cpp:2680` resolves that op to the kernel whose own header at `src/vt/cuda/cuda_ops.cu:1456-1460` calls itself "Correctness-grade (M0.9)": one 256-thread block per (query, head), a 256-wide shared-memory tree reduction for every key, and no K/V tiling. The tower has no second attention path, no knob and no A/B rung, so nothing could route around it. It now names `vt::AttentionDenseFlash`, the rung `whisper_audio.cpp:310-322` and `qwen3_vl_vision.cpp:462-480` already default to. `AttentionDenseFa2` is not usable at head_dim 96 and would fall through to the same kernel anyway. The size of the problem is now read rather than inferred. #1545 labelled its token count illustrative and asked whoever took it to pin the shipped resolution first, because it moves the estimate quadratically, and the released checkpoint was already on the NAS. `config.json` ships `layer_types` explicitly -- 13 `full_attention` and 37 `window_attention`, exactly what `muse_glimmer_weights.cpp:429-439` derives when the key is absent -- and `pos_emb_height/width` put the window at 32x32 = 1024 patches, so a window layer does not attend over the whole image. `processor_config.json` sets `max_image_tokens` to 4096. Whether that counts patch tokens or post-merge tokens is not decidable from the checkpoint, so both readings are carried: 34 s or 375 s of naive attention per image, against the 4.8 s the issue illustrated with. This lands UNREACHED, deliberately, under the staged-slice exception in `.agents/reachability.md`. What is not reached: `MuseGlimmerVisionForward` and its two callers `MuseGlimmerEncodePixelGroups` and `MuseGlimmerGenerateGreedyViaRegistry` have no caller in `src/`, in `examples/` or in `include/vllm.h`, only in `tests/`, and `muse_glimmer_registry.cpp:13-14` says why: "The perception encoder is still W3, so an image or video prompt is a pending brick." The row that owns the wiring is `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation`, at W4 and W5 of `.agents/specs/muse-glimmer.md` §3. The issue that tracks it is #1566, filed by this change because the model's umbrella #268 is closed and nothing open named the gap; both specs list it under `## Owed`. Repairing the cost before the wiring costs one line. Repairing it afterwards costs a regression hunt through a change that did not touch attention. Numerics, both halves. On CPU the swap is byte-identical BY CONSTRUCTION: `src/vt/cpu/cpu_ops.cpp:3760` registers `kAttentionDenseFlash` to the same `AttentionKernel` function pointer `kAttention` gets at `:3750`, and each op has exactly one CPU provider. Every existing golden passes at its previous value and no tolerance was widened. On CUDA it is NOT bit-identical: the flash kernel reduces per-lane head_dim strips with `__shfl_xor_sync` under a running online softmax where the naive kernel uses a 256-wide block tree, so the f32 partial-sum grouping differs, as `include/vt/ops.h:3304-3327` states. No CUDA run backs this change -- `dgx:gpu0` is held by the operator and no lease was taken -- and the A/B is recorded as owed rather than claimed. The test had to be a ROUTING test, because a numeric one cannot see this change: on CPU both ops are the same kernel, so any comparison passes before it and after it. The new case counts op-provider selections the way `tests/vllm/models/test_ltx2.cpp:665-686` does for the same question, and asserts a DERIVED 12 -- 5 window + 2 full + 5 window calls over the fixture's segmentation -- so a partial re-route fails it too. `src/vt/cuda/cuda_ops.cu` is deliberately untouched. #1544 item 2 owes a repair to `LaunchAttentionDenseFlash`'s advertised `head_dim <= 256` contract and the LTX-2.5 row is implementing it. head_dim 96 does not need it: the production bf16 arm asks 24 KiB of the 48 KiB default cap, and the spec records that the f32 arm sits exactly on the cap and fails loud rather than silently if it ever runs on CUDA. ## How to verify ```sh cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release cmake --build build --target test_muse_glimmer_vision test_muse_glimmer_wiring -j 8 ./build/tests/test_muse_glimmer_vision ./build/tests/test_muse_glimmer_wiring ``` Observed on this branch: 8/8 cases and 103 assertions, then 9/9 cases and 10,317 assertions, both `SUCCESS!`. The routing case prints `attention selections: dense-flash +12, naive +0`. Red first, with the test in place and the production line still on `vt::Attention`: `dense-flash +0, naive +12`, and both `CHECK`s fail as `CHECK( 0 == 12 )` and `CHECK( 12 == 0 )`. Byte-identity measured, not only argued. Both binaries were built from one tree with only that line differing, and every `MESSAGE:` line of the suite was captured from each and diffed. Exactly one line of fourteen differs, and it is the routing counter. The f32 stage numbers stay at `rel_l2` 1.201e-07 to 2.983e-07 and the bf16 tower at `rel_l2=5.951e-03 max_abs=3.675e-02` on both sides. `scripts/agent-preflight.sh` reports **All gates green.** on the head, run here. An earlier run of it on this branch failed `test_cpu_x86_llamacpp_floor`, and that failure was inherited rather than caused here: it failed the same way on the unmodified base `04f1cead6`, with `waiting for quiet: 15s busy=111% builders=0 load=67.47` and exit 4 where the case wants 2, which is the load-dependent behaviour #618 tracks. The host was at loadavg 67 to 121 then and lower at the final run, and the gate passed. THREE CI lanes are red on this pull request and none is caused here. Each was verified per-job against a main baseline, by comparing the failing ASSERTION and not the job name: - `windows-msvc-cpu` and `windows-msvc-vulkan` fail at `ERROR: CMakeLists.txt: MSVC /W4 /WX policy is negated on the C/C++ compile by /w`, before any compilation. They are red on main's own scheduled baseline too, and #584 / #965 track them. - `build-test-cpu` and both `sanitize-cpu` arms fail on ONE shared doctest case, `test_runner.cpp:1557`, which landed on main in `e2a9e035d` (#1273) and is owned by #1602 and #1608. The scheduled main baseline at `e2a9e035d` fails the identical `CHECK_THROWS_WITH_AS( make_runner(), "Block size must be a multiple of 16", std::invalid_argument )` with the same "No valid attention backend for device type 0" text, in all three jobs, with zero sanitizer findings in either sanitized arm. This branch does not touch `test_runner.cpp`. `build-newest-gcc` was red when this body was first written and is GREEN here: #1581 and #1618 landed the `::getpid` repair on main and this branch was re-merged onto `2e7f3bee7` to pick it up. `test_cpu_threadpool`, which an earlier run could not verify as inherited, PASSED on the re-run, leaving `test_runner` as the whole of the `build-test-cpu` red — 1 failing test out of 585. ## What is not verified No CUDA run. `dgx:gpu0` was held by the operator, no lease was taken, and the A/B that would confirm the speed is recorded under `## Owed` in the spec rather than claimed. The cost figures above are arithmetic from the 5.70 ns/iteration GB10 constant recorded in #1544, not a measurement of this tower. Closes #1545. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
localai-bot
added a commit
that referenced
this pull request
Aug 22, 2026
…de kernel, and one forward cost 47.84 s (#1557) One LTX-2.5 DiT forward at `768x448/49f` measured **47.84 s** on GB10 -- n=119, median 47.91 s, spread 5.8%, from the engine's own `last=` samples -- and first-order arithmetic says it should take well under a second. The DiT self-attention called `vt::Attention`, which on CUDA resolves to the kernel whose own header calls itself "Correctness-grade (M0.9)": one 256-thread block per (query, head), a 256-wide shared-memory tree reduction per key, and no K/V tiling, so K and V are re-read from global once per (query, head). At that geometry the video stream is 2352 tokens x 32 heads = 75,264 blocks each looping 2352 keys, over 48 layers. The attribution is arithmetic rather than assertion. `.agents/specs/multimodal-speed.md` measures that same kernel on that same box at **5.70 ns per block-key iteration**, and 1.77e8 x 5.70 ns x 48 = **48.4 s** against the measured 47.84 s. A 1% match, which leaves the other 1% for 48 blocks of GEMMs, norms, RoPE, gating and six attentions each. **After this change one forward is 7.680 s** (n=19, median, `768x448/49f`, GB10). The ratio is **~6.0x**, and it is **not** an A/B. Both qualifications are load-bearing and both are under "The A/B" below. ## Why nobody saw it, which is the part worth keeping `kAttention` is deliberately frozen on the naive kernel so that text decode stays byte-identical. That decision is correct and this change does not touch it. The consequence is the defect. The fast kernels are **separate ops that each caller must opt into by name** -- `kAttentionDenseFast`, `kAttentionDenseFlash`, `kAttentionDenseFa2`. There is no automatic selection, no shape routing, and no fallback notice. A model that never opts in gets **correct output at roughly 500x the cost, with no warning anywhere**. Nothing in this tree can detect that. The output is right, so every golden passes. The op is registered, so no refusal fires. `GetOpProviderStats` counts the naive selection and reports it as a success, because it is one. The only symptom is a wall clock, and a diffusion render has no reference wall clock to be held to. That is why the gate for this change is a dispatch observation and not a number. ## What changed **`ltx2_device.cpp` calls `vt::AttentionDenseFlash` in the self-attention branch.** That is the whole of the product change. The dispatch RULE is unchanged: the branch is still chosen by `context == nullptr && a.bias == nullptr`, upstream's own self-attention marker, never by what the numbers happen to be. Only the op it calls moves. Its square-problem contract holds by construction, because that branch is entered only when `s == tq`. **`VLLM_LTX2_DIT_FLASH_ATTN=0` restores the old op**, so both arms of the measurement run from one binary. Same shape as `VT_FA2_DENSE`, documented in `docs/ENVIRONMENT.md` as a measurement lane and never a configuration. Both remaining naive `vt::Attention` call sites -- that one, and `ltx2.cpp`'s CPU-only host arm -- carry a `// VT-ATTN-NAIVE:` line saying why, the form #1578's checker defines for a deliberate site. **The host arm at `ltx2.cpp` is deliberately NOT moved.** It computes into `std::vector<float>` and is CPU-only by construction, and on CPU both ops are the same registered function (`src/vt/cpu/cpu_ops.cpp:3750-3761`), so the swap would be a byte-identical no-op that moves the L2 parity reference off the reference op. ## The shared-memory cap-raise is REVERTED, and that is the change since the review An earlier version of this branch also raised `LaunchAttentionDenseFlash`'s dynamic shared-memory cap through `cudaFuncSetAttribute`, moved `SetDynamicSmemOptIn` onto a shared seam, and added head_dim contract cases to `tests/vt/test_ops_attention.cpp`. **All of it is gone.** `src/vt/cuda/cuda_ops.cu`, `cuda_device_caps.h`, `cuda_arch_tactics.cu`, `cuda_paged_attn.cu` and `tests/vt/test_ops_attention.cpp` are byte-identical to `main`. **The swap never needed it.** LTX renders at the stream dtype, and in production that is bf16 (`ltx2_device.cpp:1183`). The flash op's K/V tile is `2 * kFlashBc(64) * head_dim * sizeof(Tin)`, so the video stream's head_dim 128 asks **32,768 B** and the audio stream's head_dim 64 asks 16,384 B. Both are inside the **49,152 B** every CUDA architecture gives a launch without any opt-in at all. **What the raise was actually serving does not fit.** The shapes over 48 KiB here are f32: head_dim 128 at 65,536 B, and head_dim 256 at 131,072 B against GB10's queried ceiling of **101,376 B** (`cuda_device_caps.h:46`). The 256 shape therefore does not fit *even with the opt-in* -- it was falling back to the bit-identical `AttentionDenseFast` while the case reported a launch. So the raise bought this row one shape it does not run, and paid for it by moving a shared helper across two files and colliding with #1578 on the same lines. **#1578 owns the bound and takes the opposite, better approach.** Rather than raising the cap it makes the ADVERTISED domain honest: `AttentionDenseFlash` declares `head_dim <= 256` while it can only launch bf16 192 / f32 96, and #1578 narrows the declaration to what the code can do and refuses above it. That is a property of the code rather than of whichever device is underneath, and it is the `supports_head_size()` polarity vLLM already has. **#1578 merges first**, and after it bf16 head_dim 128 is inside the declared bound, so this change is unaffected. There is now nothing left to conflict. **One consequence is disclosed rather than left to be found.** With the raise gone, the f32 L2 parity arm at production geometry (head_dim 128, 65,536 B) reaches `AttentionDenseFlash` and cannot launch: a `cudaGetLastError` throw at `cuda_ops.cu:3352` today, a `VT_CHECK` naming the head_dim once #1578 lands. **It fails loud in both worlds and never silently**, and nothing gated reaches it -- production is bf16, and the f32 arm is a parity reference exercised at the fixture's reduced dimensions. Filed under `## Owed`. ## Numerics, measured rather than asserted **On CPU: byte-identical.** `kAttention` and `kAttentionDenseFlash` are the same registered function pointer, and the goldens are unmoved -- `test_ltx2_device` **22/22, 652/652**; `test_ltx2` **43/43, 4581/4581**; `test_ltx2_video` **102/102, 4194/4194**; `test_ops_attention_cross` **9/9, 32/32**. **On CUDA: NOT bit-identical, and here is the number.** The warp kernel groups the head_dim partial sums across 32 lanes instead of a 256-thread block, so the same f32 online softmax associates differently. On `dgx:gpu0` the host-vs-device parity case measures **video 8.9407e-08, audio 4.47035e-08** against its committed `2e-5`, and bf16 CUDA-vs-CPU-backend at **0**. That is f32 round-off scale, 224x inside the gate, and **the gate was not widened**. **That is the only numeric evidence there is, and it bounds less than it looks like it bounds.** The case runs the fixture's reduced dimensions, so it bounds the ARITHMETIC change -- a length-D sum reassociated -- and not the change at head_dim 128 with 2352 keys over 48 layers. A diffusion render has no token gate to fall back on, and the flash arm was interrupted before writing any frames, so **no pixel comparison exists** either, not even against the completed 49-frame baseline render already on the NAS. Filed as **#1612** and listed under `## Owed`. **The `test_ops_attention` evidence is WITHDRAWN, not restated.** The GB10 lease ran it at 10/10 and 88,439 assertions, and this change claims nothing from that run: it measured the head_dim cases that came with the cap-raise, and those are no longer in the tree. Two of its three arms would not have supported the claim anyway, for the reason above. The `test_ltx2_device` rows and the render below DO survive the revert, and that is arithmetic rather than assertion -- the measured binary carried the opt-in call, but the helper returns immediately below 49,152 B, so it was a **no-op on every launch those numbers came from**. ## Reachability, twice **Unit.** A new case drives the production entry point `Ltx2DitForwardDevice` -- called from the denoise loop at `ltx2_video.cpp:4246` -- and asserts the dispatch **two-sidedly** through `GetOpProviderStats`: `kAttentionDenseFlash` selected exactly **8** times (two self-attentions x two blocks x two batch rows) and `kAttention` selected **0**. The negative half is what makes it a routing proof rather than an addition proof. The case scope-guards its own process state, because it enables a counting instrument and sets an env var and contains `REQUIRE`s; measured with a scratch `REQUIRE(false)` and an appended observer, **without the guards the observer reads the knob still set and 8 leaked selections, and with them it reads neither**. That `:4246` is itself a repair. This branch and its spec both cited `ltx2_video.cpp:4055-4059`, which points at prose about the res_2s step counter -- at the row's declared base `6b48edb2c` and at HEAD alike, so re-reading it at either revision would have caught it. It is the citation the whole reachability argument rests on. `cpu_ops.cpp:3551-3562` was wrong the same way and pointed at `FusedStore`. Every other anchor in the spec has been re-read against the base rather than carried forward, and eight more were corrected. **Mutation M1**: restore `vt::Attention` at the call site, +1/-1, compile rc 0. Both halves went red -- `CHECK( 0 == 8 )` and `CHECK( 8 == 0 )`, exit 1 -- while every golden case in the same binary stayed green. That contrast is the finding: no numerical gate in this tree can see a 500x slower kernel that computes the right answer. Tree restored and re-gated green. **Production, on the real model.** With `VT_OP_PROVIDER_STATS=1` the full 21.00B render at `768x448/49f` on GB10 announces `op=21 device=1` (`kAttentionDenseFlash` on CUDA) and announces `op=18 device=1` (`kAttention` on CUDA) **zero times**. Same two-sided claim, taken through `--device cuda` at full scale rather than on a fixture. ## The A/B: one arm measured, the pair still PENDING Lease `6c724dfd` on `dgx:gpu0`, source `30dce3a1d`, one binary built in-lease with cutlass-nvfp4, cutlass-fp8 and FA-2 all `ENABLED for [121a]`. Correctness cleared before any speed number was read. **Flash arm**, per DiT forward at `768x448/49f` = 2352 tokens, from the engine's own `last=` lines: | n | median | mean | min | max | spread | |---|---|---|---|---|---| | 19 | **7.680 s** | 7.633 s | 7.109 s | 8.196 s | 14.2% | **The naive arm did not run, so this is not an A/B.** At forward 20 the `rc` worker was lost and `dgx:gpu0` read `unhealthy (no contact)`. The cause is UNPROVEN and this change does not name one: no memory trace was taken and the box did not return to be asked. What is established is that the harness as first written carried no memory guard and no sample cap, which is a defect in this row's harness rather than a finding about the change. **The 47.84 s denominator carries two confounds this arm does not, and both inflate the ratio.** Neither was disclosed before: - **A stack sampler.** The denominator ran under `runguard.py --stack-period 12` (`render.log:1`), which `eu-stack`s the process and so `ptrace`-stops every thread. Its own `stacks.txt` prices that: **523 samples, median inter-sample delta 12.40 s against a 12.0 s period**, so ~0.40 s median and 1.50 s max of stopped process per sample. About 3.9 samples land inside each 47.84 s forward, ~1.54 s, **~3.2%**. Correcting only the denominator gives **46.3 s / 7.680 s = 6.03x**. - **A different prompt.** `render.log:1` carries a ~70-word prompt; the harness uses one short sentence, and `ltx2_video.cpp:2253` sets `context_tokens = encoded.seq` **unpadded**, so the DiT's cross-attentions see a different number of keys in each arm. Corroborated rather than inferred: `conditioning.tower` is **45.013 s** against **28.426 s**. Same sign, and not quantified. So the defensible statement is the range **6.03x to 6.23x**, quoted as **~6.0x**, with the sampler correction named and the prompt confound uncorrected and pushing the same way. `6.23x` survives in the records only as the uncorrected upper end of that range, never on its own. The A/B gate reads `PENDING` and this change does not claim otherwise. **The flash arm's artifacts do not record what it ran**, which is why those confounds had to be established from a phase duration. `arm-flash.log` opens at `[render] + load` with no command line, `wd-flash/` is empty, no `phase-log.json` was written, and the only description of the run was a mutable NAS path edited 25 minutes after it finished. Both halves are repaired: the harness is committed as **`scripts/ltx25-dit-attn-flash-ab.sh`**, and every arm now writes its own invocation -- harness sha256, binary sha256, source SHA, geometry, seed, prompt, resolved command line -- to line 1 of its own log. It also caps each arm at 13 samples, holds a 12 GiB `MemAvailable` floor, caches the build on the source SHA, and runs the **naive arm first**. That harness also had a precondition that could never pass. It grepped `cuda_ops.cu` for `FlashTileSmemOptIn`, a spelling no revision of this change ever used, so it counted 0 and would have `exit 42`-ed on a correct tree as readily as on a wrong one. It is removed with the cap-raise it guarded. ## CI - **`documentation-checkpoint` was RED and it was THIS BRANCH's, not inherited.** `2aa78c69b` and `2f39a9426` each recorded a measurement in `.agents/benchmark-record.md` without writing `docs/STATUS.md` (and `docs/BENCHMARKS.md` for the second); the control on the main-only range `4c193bd55..5d548d0` is rc 0. Neither commit is in this branch's history any more, and the checker is **re-run at this head** rather than trusted to have stayed fixed -- a job that has stopped appearing in a failing set is not the same fact as a job that passes. Local rc 0 over the branch range. A side effect worth recording: that job runs `set -eu` and this checker is the FIRST of three commands, so **`check-now-current.py` and `check-role-discipline.py` never ran in CI on this branch at all**. Both are rc 0 locally, so nothing hides behind it. - **`build-newest-gcc` is now GREEN on `main`** since #1581 landed, and this branch carries that fix through the merge. Earlier runs of this PR predate it. A red here now would be this branch's, not inherited. - `build-test-cpu` and both `sanitize-cpu` lanes: **inherited**, and verified from this head's own logs rather than from the issue numbers. All three fail on exactly one case out of 585 -- `test_runner.cpp:1557`, `CHECK_THROWS_WITH_AS(make_runner(), "Block size must be a multiple of 16", ...)` receiving `No valid attention backend for device type 0` instead -- with byte-identical text in all three, and `main`'s own newest baseline run fails the same case. `test_ltx2_device` passes in all three, so this change's own cases are green on the lanes that run them. From #1273; owned by #1602 and #1608. - `windows-msvc-cpu` / `windows-msvc-vulkan`: **inherited**, baseline-less lane. A markdown-only control PR (#1295) fails the identical step. #584/#965 own them. - The full set: **16 pass, 5 fail**, and the five are a strict SUBSET of `main`'s newest baseline at `503e459005d7`. `scripts/main-baseline.py` was the instrument, not the push runs, which are all cancelled (#274). ## Owed, filed and not folded in - **#1612** -- there is no numeric or pixel comparison at production geometry. The only numeric gate is the reduced-dimension one above; the flash arm wrote no frames, so no pixel A/B exists. The f32 parity arm's refusal at production geometry is recorded against the same issue. - **#1551** -- `vt::AttentionDenseFa2` still refuses head_dim 128, so LTX cannot reach tensor cores. Everything here is still a scalar warp-per-query recurrence. - **#1552** -- the same opt-in-by-name defect reaches every other `vt::Attention` caller, and nobody will be told there either. All three are listed under `## Owed` in `.agents/specs/ltx25-dit-attn-flash.md`, together with the two `scripts/attention-rung-allowlist.txt` stems that #1578's checker will report `STALE` once the markers here meet it. Closes #1549 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
localai-bot
pushed a commit
that referenced
this pull request
Aug 22, 2026
…h review falsified The merge first, because it is the hazard this branch's own #1619 row documents. `origin/main` moved to `5539686c7` (#1672) during review and it appended a row to `.agents/issue-index.md`. Resolved by the canonical rule rather than by trusting the union driver: 586 rows, 586 unique ids, and `git diff --numstat origin/main -- .agents/issue-index.md` is `10 0`, additions-only against the moved main. Then the part that matters. This is a record pull request, so a false sentence in it is the whole defect and not a blemish -- and four of its claims came verbatim from #1556's spec without being re-verified. A record change that propagates a superseded change's unchecked assertions repeats the failure it exists to prevent, one remove further out. Two of them were landing in append-only rows. **`check-agent-record.py` never passed the duplicate.** #1556's spec said `git merge-tree`, `check-agent-record.py` and `check-issue-index-append-only.py` all accepted the 538-line union output. Reproduced here, and it is false: that checker returns rc=1 with `issue #1546 listed twice`, and the refusal has existed since `8dd6508da` (2026-08-09), before the merge in question. Verified twice -- once by regenerating the driver output at the historic merge, once by duplicating the last row on this head -- and restored byte-for-byte both times. The gap is real and NARROWER: exactly one checker is blind, not two. **"Three months" is arithmetically impossible.** The four issues that argued the tolerance -- #1439, #1470, #1494, #1536 -- were filed between 2026-08-20T05:23Z and 2026-08-21T04:40Z, under 24 hours. This repository's first commit is `accfae2de`, 2026-07-02, so nothing in it can be three months old. **`ltx2_res2s.cpp` has never existed.** `git log --all --diff-filter=A` finds no such file. `Ltx2Res2sDenoisingLoop` lives in `ltx2_samplers.cpp`, is declared in `ltx2_samplers.h` beside `Ltx2Res2sHooks`, and is called from `ltx2_video.cpp`. #1567's own forge text names no file, so the wrong anchor was invented by the spec and would have been copied forward by anyone who read the row. **The `#838` recurrence is withdrawn as unreproducible.** Re-running `git merge-file --union` at every later merge where `#838` sat on a side leaves it present in all of them. A union driver cannot drop a line that is an addition on one side; a wholesale take-ours resolution can, which indicts the resolver. That sentence was what turned a measured one-off into a recurring class, so it carried weight it had not earned. Two smaller ones. `d27639e71` is `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` (#1476/#1595) and added the same `<unistd.h>` in passing; only `13548db8f` (#1581) is the `build-newest-gcc` repair. And "records that as a holding action, in terms" claimed a quotation the tree does not contain -- the substance is right, the words are not, so it now says so. Also stated rather than implied: this row has no matrix row and therefore no lifecycle state, which is deliberate while the implementation is unlanded, and #1668 is the forward owner that creates one. And "neither red reproduces" is scoped to the two floors this row was filed against, because #1572 in this same batch records assertion (1c) redding intermittently on `main`. Preflight exits 0 with all gates green and none skipped. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot
added a commit
that referenced
this pull request
Aug 22, 2026
… request, and ten open issues had no index row (#1670) PR #1556 is being closed rather than merged, and this lands the record it would otherwise have taken with it. ## Why that pull request cannot merge `squash_merge_commit_message = PR_BODY`, so its body becomes the permanent commit message. That body argues at length for a `load.setup` anchor that `519303d15` (#1622, row `LTX25-DEVICE-RESIDENCY`) had already landed as `load.open` — same region, same open and close statements, same `Scope::Close` shape. Merging it writes a materially false narrative onto `main` irreversibly. **That is a defect in the message, not in the code.** ## Why the code alone was not enough to salvage Verified on `db648fb88` by a tree-wide grep over `src/`, `include/`, `tests/` and `docs/`: | symbol | occurrences on `main` | |---|---:| | `load.dit_config` | 0 | | `artifacts.mux` | 0 | | `sampler_updates` | 0 | | `Record::instrument_seconds` | 0 | | `denoise.update` | 1 — `test_ltx2_video.cpp:4325`, which is `6b48edb2c`'s comment ASKING for it | So the supersession covers ONE anchor, not the branch: three anchors plus a ~193-line change to shared instrument semantics and 583 lines of test remain genuinely new. That is a fresh row's work with its own red-first tests and its own fresh review, on a branch that conflicts in three files — not a small salvage. It is owed under #1668. ## What this commit preserves Ten issues are open on the forge and **none of them had a row in `.agents/issue-index.md`**: #1470, #1536, #1567, #1568, #1569, #1570, #1571, #1572, #1619 and #1668. Closing the pull request without this would have left all ten invisible in the record, along with 809 lines of measurement. Two things in there are worth more than the anchors: **The residue decomposition.** 92.09% of the un-named time is ONE region — the load prologue, 17.661 ms of a 19.178 ms residue. The other four gaps are 4.95%, 1.30%, 1.09% and 0.56%, and the sixteen gaps between adjacent named phases hold 6.8 us each, which is the instrument and nothing else. #1536's own first hypothesis, `d995c52f0`'s temporal x2 upsampler, is **refuted** rather than deprioritised: it runs inside `phase.upsample_latent`, a named leaf, and is not in the residue at all. **A rejected gate design, with its distributions.** `residue <= 2 * instrument` measured red 4 in 45 at the table bound (load 88, max 4.115), 3 in 200 at the conservation case, 2 in 200 at `unit.parent`, and 28 in 160 (17.5%, max 5.55) on a standalone probe at load 125. The un-instrumented remainder of a boundary dilates faster than the instrumented part under contention, so the ratio has a heavy right tail — and the 20-run distribution that read 1.021 to 1.464 saw none of it. Three fresh reviews paid for that. `## Design` 3 exists so nobody measures it again. ## The spec is written to today's position No `## Gates` section, because the row is not runnable and a `## Gates` edit moves a row into the runnable population as a side effect. No `## Outcome`, because there is no outcome yet. `## Now` states plainly that the implementation is not on `main`, and names `refs/pull/1556/head` = `b45ea3bbb`, where the reference implementation, the full gate report, the mutation table and the three review threads stay readable. Two index rows are corrected **before** they are appended rather than after. The branch wrote #1470 and #1536 as closed by anchoring the sampler step — true of the branch, never true of the tree. They now say what actually landed instead, and say that an earlier draft claimed otherwise. ## Five inherited claims the fresh review falsified, and this record does not repeat A record pull request whose sentences are wrong is entirely defective, and four of these came verbatim from #1556's spec without being re-verified. Two were landing in append-only rows. Repaired in `2f2fc2dc7`. **`check-agent-record.py` never passed the duplicate.** #1556's spec said `git merge-tree`, `check-agent-record.py` and `check-issue-index-append-only.py` all accepted the 538-line union output. Reproduced, and false: that checker returns **rc=1** with `issue #1546 listed twice`, and the refusal has existed since `8dd6508da` (2026-08-09) — before the merge in question. Verified twice, restored byte-for-byte both times. The gap is real and **narrower**: one checker is blind, not two. **"Three months" is arithmetically impossible.** #1439, #1470, #1494 and #1536 were filed between `2026-08-20T05:23Z` and `2026-08-21T04:40Z` — under 24 hours. This repository's first commit is `accfae2de`, 2026-07-02, so nothing in it can be three months old. **`ltx2_res2s.cpp` has never existed.** `git log --all --diff-filter=A` finds no such file. `Ltx2Res2sDenoisingLoop` lives in `ltx2_samplers.cpp`, declared in `ltx2_samplers.h` beside `Ltx2Res2sHooks`, called from `ltx2_video.cpp`. #1567's forge text names no file, so the wrong anchor was invented by the spec. **The `#838` recurrence is withdrawn as unreproducible.** Re-running `git merge-file --union` at every later merge where `#838` sat on a side leaves it present in all of them. A union driver cannot drop a line that is an addition on one side. That sentence was what turned a measured one-off into a recurring class. **`d27639e71` is not a `build-newest-gcc` repair.** It is `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` (#1476/#1595) and added the same `<unistd.h>` in passing; only `13548db8f` (#1581) is that lane's fix. And "records that as a holding action, in terms" claimed a quotation the tree does not contain — right in substance, not in those words. ## A fifth claim — this one was mine, not inherited While waiting on CI I re-measured the one assertion in this PR I had taken from the reviewer rather than run myself. It was overstated, and it was landing in an append-only row. Repaired in `72589fe6e`. The #1619 row said `check-issue-index-append-only.py` reds a correct de-duplication. Measured on a scratch branch off this head, **with commits** — that checker diffs `merge_base..HEAD` from Git objects and ignores a working tree: | committed mutation | `check-issue-index-append-only.py` | `check-agent-record.py` | |---|---|---| | duplicate a row | **rc=0** (blind) | **rc=1**, `listed twice` | | then remove the duplicate | **rc=0** | — | | relocate a **base-reachable** row | **rc=1**, `removed: ... #168` | — | Add-then-remove nets to nothing against the merge base, so it passes. **And the de-duplication claim needs a condition, which a re-review supplied and which I had not measured.** The rows above all use `origin/main` as the base — a base that *predates* the duplicate. Run at three pairings: ``` --base e2a9e03 --head <real canonical 537> -> OK rc=0 --base e2a9e03 --head <synthetic 537> -> OK rc=0 --base <committed 538> --head <537 de-dup> -> FAIL rc=1 ``` So the checker reds a de-duplication **only when the duplicate is already reachable from the base** — when you are repairing a corruption that already landed. It does not fire on someone who fixes union output before committing it, which is the ordinary case and the shape this branch used. The same range property is what makes relocating a base-reachable row red. Two lessons recorded rather than absorbed silently. Mine: three earlier probes edited the **working tree** and read rc=0 from all of them, including one that deleted a base row — that is not evidence about a checker that reads commits. The reviewer's, which it volunteered: it first marked this half confirmed from **reading** `removed = [...]`, and reading the collector says what it collects, not what the diff *range* puts in front of it. Recorded because the method matters: my first three probes edited the working tree and read rc=0 from all of them, **including one that deleted a base row**. That is not evidence about the checker, it is a restatement of the fact that it reads commits. The `check-agent-record.py` rc=1 result re-confirms the F1 correction above in the same run. ## Evidence ``` git diff --numstat <merge-base> HEAD 10 0 .agents/issue-index.md .agents/specs/ltx25-phase-residue.md (new file) issue-index: 592 rows, 592 unique issue ids, 0 duplicates, additions-only this branch's ten rows are 598-607 (#1470 .. #1668); main's #1666/#1667 follow at 608-609 and #1131 at 610, so the file ends at neither side's append scripts/agent-preflight.sh exit=0, All gates green, none skipped ``` The evidence quotes a **merge-base sha, not `origin/main`** — deliberately. `main` moved three times during this work (`db648fb88` → `5539686c7` → `9712ab77f` → `3e5d2f370`), and the same command against a moved main measures a different thing: `git diff --numstat 5539686 3d6a0a4` reads **`10 1`**, where the lone deletion is main's own `#1131` row that the pre-merge head did not yet have. Against the merge base it reads `10 0`. A number derived from a name that resolves differently every few minutes describes whichever main was current when the command ran, not the change. (An earlier revision of this body said that reading was `1 1`. It is not reproducible at any pairing, and it is corrected here rather than carried — the index did not move at all between `5539686c7` and `9712ab77f`.) **On the row placement, stated rather than glossed:** the canonical procedure the #1619 row prescribes is to append at the true end and re-verify. That is not what this head looks like — the union merge placed main's `#1131` after this branch's ten. Additions-only and duplicate-free both hold, and no row was lost or edited, so the result is correct; but the ordering is the driver's, not the procedure's, and saying otherwise would repeat exactly the kind of unchecked claim this PR exists to correct. Zero-duplicate ids is asserted rather than assumed. It is **not** the detector #1619 was filed about — that was this PR's own earlier position, and the repair overturned it: `check-agent-record.py` has refused duplicate ids since `8dd6508da`. #1619's remaining ask is row **multiplicity**. `test_cpu_x86_llamacpp_floor` red once during this work at loadavg 31.79, `NO_QUIET_WINDOW` (4) where the case expects `GIVING_UP` (2). That is #618 verbatim and load-dependent: it passed on a clean tree and then passed again on this tree at lower load. A records-only diff of two Markdown files cannot reach a CPU floor harness. One scope correction while we are here: "neither red reproduces" means the two floors this row was filed against. It is **not** a claim that `test_ltx2_video` is quiet — #1572, one of the rows in this batch, records assertion (1c)'s span slack redding intermittently on `main`, which this row neither causes nor repairs. This row has no matrix row and therefore no lifecycle state, so `scripts/now.py` and `audit-live-rows` will not surface it. That is deliberate while the implementation is unlanded: creating a row now would put an empty one in the runnable population. #1668 is the forward owner and creates the row with the code. ## Nothing is closed here #1439, #1470, #1536, #1567, #1568, #1569, #1570, #1571, #1572 and #1619 all stay open. #1668 stays open and carries the implementation. This commit records #1668 rather than discharging it. ## Why this merged without a complete rollup on its exact SHA Recorded because a reader six months from now will see a merge on a partially-settled rollup and needs to know it was a judgement, not an oversight. `main` moved five times while this branch was in review (`db648fb88` → `5539686c7` → `9712ab77f` → `3e5d2f370` → `d25730fbb` → `4ede6bdfd`). GitHub does not honour the `merge=union` driver on `.agents/issue-index.md`, so each move marked this PR `CONFLICTING` and forced a real merge commit, which restarts a ~75-minute CI cycle. At ~10 commits/hour into main, that loop has no fixed point. **The gated tree is `d25730fbb`** — a commit on `main`, carrying `main`'s own CI verdict. This branch's head `b3761b976` is that tree plus this change, and the delta cannot reach anything CI compiles or runs: ``` git diff --stat d25730f b3761b9 -- src include tests examples [empty] git diff --stat d25730f b3761b9 .agents/issue-index.md | 10 ++ .agents/specs/ltx25-phase-residue.md | 285 +++++++++++++++++++++++++++++++++++ 2 files changed, 295 insertions(+) ``` Two Markdown files, **295 insertions, zero deletions**. The compiled artifacts at this head are byte-identical to `d25730fbb`'s by construction, so `d25730fbb`'s verdict is not "probably still valid" — it is a verdict about the same binaries. Both SHAs are pinned deliberately: `origin/main` is a moving ref, and diffing it instead of a pinned commit made this same branch appear to delete 768 lines of `vocoder1d.cpp`, `cpu_conv1d_*` and three test files, when the deletions were `4ede6bdfd` landing between two commands. What did settle on this branch's own rollups was verified by failure **text**, not job name, against `main`'s baseline `90e8c3c85`: `build-test-cpu` and `sanitize-cpu (address,undefined)` each fail one test of 590, test #282, on a **byte-identical 335-byte** assertion at `test_runner.cpp:1557`, with **zero** sanitizer diagnostics on both sides (#1602/#1608); `windows-msvc-cpu` and `-vulkan` fail on the `/W4 /WX ... negated by /w` checker refusal (#1649) that fires before any compilation, with zero `error C####` and zero `error LNK####`. `build-newest-gcc` — the one job that can never be called inherited, since `main-baseline.py`'s newest row predates #1581 — is **green**. Records verified on the merged result rather than against the moving ref: 594 rows, 594 unique issue ids, zero duplicates, and `git diff --numstat` additions-only for this branch's ten rows. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
localai-bot
added a commit
that referenced
this pull request
Aug 22, 2026
…he file it named (#1665) `47a918d8f` (#1579, issue #1545) and `90e8c3c85` (#1557, issue #1549) landed the routing that `scripts/attention-rung-allowlist.txt` 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 whoever runs preflight next. This is that preflight. Closes #1663. ## The window is not free, and that is the finding `## Risks/decisions` D4 of the spec says the deferral is safe because a stale entry is reported and never fatal. True of the exit code, and not the whole story: an allowlisted stem excuses its ENTIRE translation unit, never only the call that earned the entry. Every `VT-ATTN-NAIVE:` marker in a covered file is therefore decorative for as long as the stem sits there. Measured on `db648fb88` by deleting the live marker at `src/vllm/model_executor/models/ltx2.cpp:959`: | allowlist | checker | report | |---|---|---| | the three stems, as `main` has them | **rc=0** | `7 carry a recorded reason, 1 unmarked and excused` | | the stems removed, as here | **rc=1** | `src/vllm/model_executor/models/ltx2.cpp:966` | Both arms restored byte-for-byte against a sha256 taken before the mutation. So the entries were a live hole in the guard #1544 exists to be, open from `90e8c3c85` until here — not untidiness. The deferral D4 designs for is still right, because it genuinely keeps the removing row off a shared file; what is added to the spec is what it costs, so the next row that parks a stem reads it with the cost attached. ## The three earned their green separately, and asymmetrically - `muse_glimmer_vision` names `vt::Attention` **nowhere**: `47a918d8f` routed the perception encoder's sole path to `vt::AttentionDenseFlash`. - `ltx2` and `ltx2_device` still **name** it, at calls that now record their own reason — the host arm, CPU-only by construction, where `kAttention` and `kAttentionDenseFlash` resolve to the same registered function (`src/vt/cpu/cpu_ops.cpp:3750-3761`); and the `VLLM_LTX2_DIT_FLASH_ATTN=0` arm of a same-binary A/B, which exists so both halves of the 47.84 s / 7.680 s measurement run from one build. One assertion covering all three would be false of one of them in either direction, so the new test case asserts them apart. ## The expected set moves in the same change The allowlist's own header requires it: the checker does not fail on a stem set that has drifted, and `tests/scripts/test_check_attention_rung_consistency.py::ShippedTreeTests::test_allowlist_holds_only_the_in_flight_stems` does, on an addition or a deletion alike. An empty expected set is not a weaker assertion. `drift_sites` now excuses nothing, so `test_shipped_tree_is_green` measures the shipped tree on its markers alone, and a silent append still reds this case exactly as before — verified by appending `some_new_tower` and measuring 3 failures, restored after. The new `test_the_formerly_allowlisted_stems_pass_on_their_own_merit` re-states positively what the allowlist used to assert by omission. The header keeps the reason each stem left, so an empty parking lot does not read as an abandoned one. ## Evidence ``` python3 scripts/check-attention-rung-consistency.py OK (attention rung): 8 vt::Attention call site(s) in 8 model source file(s); 8 carry a recorded reason, 0 unmarked and excused by 0 allowlisted in-flight stem(s). rc=0 python3 -m unittest tests.scripts.test_check_attention_rung_consistency Ran 35 tests in 5.297s -- OK scripts/agent-preflight.sh exit=0, All gates green. scripts/agent-preflight.sh --staged exit=0, All gates green. ``` `.agents/specs/attention-rung-visibility.md` `## Now` records the discharge and the measured cost of the window. #1629 and #1631 are separate defects in the same checker and are untouched. ## Fresh review, and the four findings repaired in `f64effeca` The review confirmed the headline mutation independently and found four things. **F1 was the one that mattered**, and it was a regression this branch introduced. The `#1631` owed entry was edited to say the checker comment at `:252-255` "denies an equality that held ... and no longer does" — wrong in the direction that misleads, because if the equality no longer held then 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 it cannot fail. The same sentence credited the 8/8/0 triple to #1663, when `db648fb88` already read 8 sites and 8 marked — `47a918d8f` and `90e8c3c85` moved that; #1663 moves only the excused count. **F2** — the marker is at `ltx2.cpp:959`, not `:958`, which is `a.causal = false;`. Corrected here 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 read `checker green (9 sites / 6 marked / 3 unmarked and excused by the 3 in-flight allowlisted stems)`. Every number was already false on `db648fb88` and all three would be false again after this row, 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, because restating it rebuilds the drift lock AGENTS.md `## Records` forbids — and that is the same defect #1629 names in this checker's own test. **F4** — `test_every_allowlisted_stem_names_a_real_model_source` now iterates over an empty set. Dormant, not dead, and the file now says so: it guards a file that is currently empty and fires on the first thing added to it, measured rather than asserted, since appending one bogus stem reds it together with the pinning case. `test_cpu_x86_llamacpp_floor` red twice during this work, at loadavg 31 with `NO_QUIET_WINDOW` and `busy=161%`. That is #618, and it is not reachable from a diff of three Markdown files and one Python comment: the same suite passed on a clean tree and on this tree at lower load. ## Why this merged without a complete rollup on its exact SHA `main` moved six times while this branch was in review, the last move being this session's own #1670 landing ten rows into `.agents/issue-index.md`. GitHub does not honour the `merge=union` driver on that file, so each move marked this PR `CONFLICTING` and forced a merge commit, restarting a ~75-minute CI cycle against a branch that merges ~10 commits/hour. That loop has no fixed point. **A complete rollup exists, and it covers every byte this branch contributes.** At `22ee35d7e` the rollup settled with **zero pending**. Against that head, the four files this branch owns are byte-identical to what is being merged now: ``` git diff --stat 22ee35d b974837 -- \ scripts/attention-rung-allowlist.txt \ tests/scripts/test_check_attention_rung_consistency.py \ .agents/specs/attention-rung-visibility.md \ .agents/kernel-matrix.md [empty] ``` So the `+41/-1` in the Python test — the one real test change here — was **present and gated** at the head that produced the complete rollup. Everything added to the range since is `main`'s own work arriving through three merges, each commit of which carries `main`'s own verdict. This is not "the delta is small enough to ignore"; it is that the delta contains none of this change. Both SHAs are pinned deliberately. `origin/main` is a moving ref, and diffing it instead of a pinned commit made a sibling branch appear to delete 768 lines of `vocoder1d.cpp`, `cpu_conv1d_*` and three test files, when the deletions were another commit landing between two commands. **Every failure at that gated head was verified inherited by failure text**, not by job name, against `main`'s baseline `90e8c3c85`: | job | evidence | |---|---| | `build-test-cpu` | 1 of 590, test #282, **byte-identical 335-byte** assertion at `test_runner.cpp:1557` | | `sanitize-cpu (address,undefined)` | same assertion, **0** ASan/UBSan diagnostics both sides | | `sanitize-cpu (thread)` | same assertion, **0** TSan diagnostics both sides | | `windows-msvc-cpu` / `-vulkan` | `/W4 /WX ... negated by /w` refusal (#1649), fires before compilation, **0** `error C####`, **0** `error LNK####` | | `build-newest-gcc` | **green** — the one job that can never be called inherited, since `main-baseline.py`'s newest row predates #1581 | The `test_runner.cpp:1557` failure is #1602/#1608: the attention-backend selector refuses a non-multiple-of-16 block size before `CheckKvCacheShape` can, so the case asserts a message only a ROCm/FLASH_ATTN build produces. Nothing outside that characterised set appeared on any run of this branch. Records verified on the merged result: 598 rows, 598 unique issue ids, zero duplicates, `1 0` for this branch's single #1663 row. Local preflight green at the pushed head, no gate skipped. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
build-newest-gcchas been RED onmainsince5702d8f83, and the finding isnot a missing include. It is a recurrence of the exact class
tests/support/process_id.hwas created to close.Closes #1575.
The seam exists, and this test bypassed it
tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547calls::getpid()whileincluding nothing that declares it. The lane builds in a
gcc:16container,where libstdc++ no longer supplies the transitive
<unistd.h>, so the file doesnot COMPILE:
Every local toolchain here is gcc 13 and still supplies that include, which is
why the lane is the only place it shows and why a locally-green preflight said
nothing about it.
The seam predicted this in its own header:
W5's GGUF test copied the temp-directory idiom and the direct call with it. That
is the copy-forward the seam was built to stop, happening again.
The convention is otherwise unanimous. Every other file in the tree naming
getpideither includes<unistd.h>or uses the helper, and the siblingtests/vllm/entrypoints/test_dspark_draft_routing.cppuses the helper ANDcarries a comment warning not to reintroduce this class. This file is the only
one that does neither.
So the repair is the seam, not
<unistd.h>. The POSIX spelling would turn thelane green and re-copy the idiom the helper exists to centralise, and it does
not compile on MSVC at all.
Attribution
Read off the baseline lane rather than inferred:
build-newest-gccis absentfrom the failed list at
92406c620and present at5702d8f83, the commit thatadded the file (SPEC-DFLASH2 W5, the GGUF drafter arm).
Evidence
Measured in the lane's own
gcc:16container,-fsyntax-onlyon the realcompile command from
compile_commands.json, at base947e5f648, file sha25683bba319051c654174fe3b3b5af72f304fdd89025443aac17ca5db91e3cc0a89:gcc:16RED_RC=1, error byte-identical to CIgcc:16GREEN_RC=0,compile_err=0GCC13_RC=0,compile_err=0The red arm was produced by stashing the fix on this branch and re-verifying the
file's sha256 back to
83bba319…first, so the red is the base bytes and not adifferent file.
git diff --statwas read on the green arm as well, so amutation that silently failed to apply could not read as a pass.
Gate
scripts/agent-preflight.sh --fail-on-skip: ZERO skips,RANGE_COUNT=2(one fix commit, one merge),
commit-trailersok andcommit-styleok,both RUN against
origin/main5d548d003.Those two gates SKIPPED on the first attempt, with the reason "Neither gate
reported anything about this tree", because
origin/mainhad moved and was nolonger an ancestor of HEAD. The merge commit here exists to make them run, and
it carries the trailer block itself rather than being a
--no-editmerge.One gate fails and it is inherited, not this change:
test_cpu_x86_llamacpp_floor. Itstest_a_contended_leg_is_discarded_and_never_summarisedexpects exit 2(
GIVING_UPafter 24 discards) and gets exit 4, whichscripts/cpu-x86-llamacpp-floor.sh:127emits asNO_QUIET_WINDOW. The host wasat load 86-118 throughout, so the quiet-window wait trips before the discard path
is ever reached. The harness is behaving correctly; the self-test is what assumes
a quiet host.
agent-record-- the CI job that runs this same file atci.yml:271-- is green onmainin run32464632413, and
this branch touches none of it. Same class as #433.
Scope
Three files: the two-line seam fix, one appended
.agents/issue-index.mdrow,and a link in the DFlash2 spec header. No production code, no behaviour change,
no other test touched.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]