fix(MUSE-GLIMMER-VISION-ATTN-FLASH): route the perception encoder off the correctness-grade attention kernel (#1545) by localai-bot · Pull Request #1579 · mudler/vllm.cpp

localai-bot · 2026-08-21T10:14:52Z

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

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 CHECKs 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 test_openai_api_server crashes on Windows with STATUS_STACK_BUFFER_OVERRUN (0xC0000409), unmasked by the #512 fix #584 / windows-msvc-cpu/vulkan are red on EVERY pull request: C4456 'loaded' shadow in server_main.cpp, and it is not #645 #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 (test(BACKEND-ROCM): the block-size contract is enforced at its production call site #1273) and is owned
by test_runner's new block-size case is red on a default CPU build: the SELECTOR refuses before CheckKvCacheShape can, so it asserts a message only a ROCm/FLASH_ATTN build produces #1602 and test_runner "refuses a non-multiple-of-16 block size" is RED on main at e2a9e035d for every CPU build: the attention registry refuses first, with another message #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]

… the correctness-grade attention kernel (#1545) `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. The size of the problem is read rather than inferred. The released checkpoint was already on the NAS, so `config.json` gives the 13 full and 37 window layer types directly, `pos_emb_height/width` put the window at 32x32 = 1024 patches, and `processor_config.json` sets `max_image_tokens` to 4096. That is 34 s or 375 s of naive attention per image depending on whether the count is patch or post-merge tokens, against the 4.8 s #1545 illustrated with. This lands UNREACHED, under the staged-slice exception in `.agents/reachability.md`. `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 row that owns the wiring is `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` at W4 and W5, the issue that tracks it is #1566, and both specs list it under `## Owed`. On CPU the swap is byte-identical by construction, because `cpu_ops.cpp:3760` registers `kAttentionDenseFlash` to the same `AttentionKernel` pointer `kAttention` gets at `:3750`. On CUDA it is not bit-identical: the flash kernel groups its f32 partial sums per lane where the naive kernel uses a 256-wide block tree. No CUDA run backs this change, and the A/B is recorded as owed. The new test is a routing test, because a numeric one cannot see the change on CPU. It counts op-provider selections and asserts a derived 12. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…nder the branch `origin/main` moved twice while this row ran, so the branch takes those commits rather than pushing a non-fast-forward. `.agents/issue-index.md` was the only shared file both sides touched, and both sides only appended, which is what the `merge=union` driver exists for. The merged file carries all six new rows, this row's two among them, and no existing row changed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…review checked and one owed item it found missing (#1545) The review of #1579 checked every anchor in the spec instead of reading past them, and four did not name what the text says they name. None of them changes a conclusion; each of them costs the next reader the trust that the rest are right. The worst was a claim of absence. The spec said there is no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`, and there are seven files that call it, one of them inside the `flash_attn/` subtree this very kernel was ported from. The claim the argument needs is narrower and true: `cuda_ops.cu` contains none, so `LaunchAttentionDenseFlash` never opts in and the 48 KiB default cap is the real one. The sentence now says that, and names the seven so a reader who greps does not conclude the argument is broken. The other three: the bit-identity quote is at `ops.h:3313-3314` and the spec cited `:3320-3327`, which is the `AttentionDenseFa2` comment stating the opposite; the 5.70 ns constant is stated in #1544 and derived from `multimodal-speed.md:24-26` rather than recorded there, and that file does not contain the number; and the test comment cited `cuda_ops.cu:1456-1459` for a sentence on line 1460, disagreeing with the production comment beside it. The spec also now owes the f32 launch probe. Section 6 argued the f32 arm fits the 48 KiB cap to the byte instead of measuring it, and an inference at an exact boundary that nobody has run is debt, not a settled result. It is safe today only because that arm is dead twice over and fails loud, which is a reason to record it rather than a reason to stop writing it down. Section 6 gains one check the review made and the spec had not: the kernel body declares no static `__shared__`, so nothing competes with the dynamic request. One static byte would put the f32 arm over. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…nder the branch The fresh review of #1579 noted the head had fallen one commit behind `origin/main` again, which makes `check-commit-style.py --range origin/main..HEAD` refuse with `range base must be an ancestor of range head` for a reason that has nothing to do with the commits. The branch takes the commit rather than pushing a non-fast-forward. `.agents/issue-index.md` union-merged as before, both sides appended only, and no existing row changed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…venance asserted an absence the tree contradicts (#1545) The previous commit fixed a false claim of absence and then made a smaller one in its replacement. It said the second anchor for the 5.70 ns constant lives in [#1544](#1544) and "not in this tree". The head_dim provenance is indeed only in the issue. The anchor is not: `.agents/benchmark-record.md:4654` records the Voxtral Whisper encoder forward over 1500 frames and 32 layers at 8870 ms, which is the measurement the anchor compares against. Only the 8.21 s prediction and the "within 7%" comparison are the issue's alone, and the sentence now says that. This is the same shape as the claim three lines above it, which is why it is worth its own commit body: a sentence that says where something is NOT costs nothing to write and takes one grep to falsify, and a reader who falsifies one stops trusting the rest. Found by the fresh review of the repair. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…nder the branch `origin/main` moved again during the repair round. The branch takes the commit so the range gates have an ancestor to work from and the push stays a fast-forward. The incoming change touches `tests/vllm/v1/worker/test_runner.cpp` only, which this branch does not, so there is nothing to reconcile. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

localai-bot · 2026-08-21T17:11:16Z

Review record, and the owed clause verified on the branch

Posted by the landing session so this PR stops showing zero reviews. Provenance
first: I am not the fresh reviewer. The fresh scoped review of this branch
returned PASS and was relayed to the operator verbally, so it never reached
this page. What follows is that verdict recorded, plus what this session
checked itself. Nothing below is a second review.

The owed clause fix has already landed here

The review's one owed repair was at .agents/specs/muse-glimmer-vision-attn-flash.md:151-152,
which claimed the second anchor of #1544
is "not in this tree". It is. Verified at this head:

The spec now reads: "The tree holds the second anchor's underlying measurement
at .agents/benchmark-record.md:4654 — the Voxtral Whisper encoder forward
over 1500 frames and 32 layers at 8870 ms — and only the 8.21 s prediction and
the 'within 7%' comparison are the issue's alone."
.agents/benchmark-record.md:4654 on this branch is the table row
| Voxtral Whisper encoder forward (1500 frames, 32 layers) | **8870 ms** (8858-8882) | ...,
read at the branch head rather than in the working tree.

So the clause is repaired and it is repaired accurately: the measurement is in
the tree, and only the prediction and the comparison are the issue's alone.

CI, read per job against a `main` baseline

main is RED. scripts/main-baseline.py at e2a9e035dbf8 reports
build-test-cpu, build-newest-gcc, sanitize-cpu (thread),
sanitize-cpu (address,undefined), windows-msvc-cpu and windows-msvc-vulkan
failing, and windows-msvc-* failing on all ten of the last ten baselines.
Read per job rather than off a rollup, because sanitize-cpu is
continue-on-error and a run's own conclusion can say success while it is red
(#274).

job on this PR	verdict	attributed
`build-newest-gcc`	FAILURE	inherited. `test_qwen3_dflash2_gguf.cpp:547: error: '::getpid' has not been declared` — byte-identical to the baseline lane, and the subject of #1565/#1575. Cleared by #1581.
`sanitize-cpu (thread)`	FAILURE	inherited. Sole failing test `280 - test_runner`, the same one the baseline's own thread lane fails.
`sanitize-cpu (address,undefined)`	FAILURE	inherited. Sole failing test `280 - test_runner`, same as baseline.
`build-test-cpu`	FAILURE	partly inherited. `280 - test_runner` is the baseline's failure. `456 - test_cpu_threadpool` is not: it PASSED in 0.34 s on the baseline run and failed here on `CHECK( ratio < 100.0 )` at `128.28` — the empty-op dispatch timing guard, a runner-contention flake rather than anything this diff can reach. Not accepted as inherited on that evidence; a fresh run settles it.
`windows-msvc-cpu` / `windows-msvc-vulkan`	FAILURE	inherited. Red on every one of the last ten `main` baselines.

test_runner is a live main red owned elsewhere and already filed twice, as
#1602 and
#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. It arrived
with e2a9e035d and no PR in this campaign owns it.

python3 scripts/agent-pr-body.py --pr 1579 exits 0 — the body will land
clean as the squashed commit message. Its only closing keyword is Closes #1545
at line 106, and no stray keyword elsewhere in the prose closes an issue this
change does not fix.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]

localai-bot · 2026-08-21T17:20:04Z

Fresh review: PASS

Recording a fresh scoped review that was run and reported verbally but never
posted, so the record existed nowhere on this pull request. Relayed here by the
operator, with the two transfer facts re-derived rather than taken on report.

Verdict: PASS.
Scope reviewed: a9c2317d8a9b427e5e06b83232ae8e75d610afda.

The reviewer inspected the change statically and mutated each claimed guarantee
in a scratch copy, restoring the tree after each. It verified findings F1 to F5
and F7, and points 7 and 8, by re-derivation rather than by reading the report of
the previous round.

The load-bearing part is why the earlier round's mutation evidence still applies
to the reviewed head. The routing change and its mutation proofs were established
at 54e1b1183, and two things carry them forward unchanged:

git diff --name-only 32a7a2f86..54e1b1183 -- src/ is empty. The repair
commit touched no product source at all.
sha256(src/vllm/model_executor/models/muse_glimmer_vision.cpp) is unchanged at
72a611ecfe857fdc3ff7e14dd0374a74be7cc971567968efdb990bdf1f08f2fd.

So the repair moved citations and records only, and the mutation evidence for the
vt::Attention to dense-flash routing transfers to that head byte-for-byte.

Operator re-derivation

Both facts reproduce, and the second one reaches further than the reviewed scope.
sha256 of that file is still 72a611ec...f2fd at the current head
92712a0b022f101ec51f7d5b78b86863f1eefada, so the mutation evidence transfers to
the head that is being merged, not only to the head that was reviewed.

Two commits landed after the reviewed scope. e6f0f01b2 is the repair of the one
owed finding this review raised: the spec asserted that the second anchor for the
5.70 ns/iteration constant was "not in this tree", which the tree contradicts.
.agents/benchmark-record.md:4654 holds the Voxtral Whisper encoder forward over
1500 frames and 32 layers at 8870 ms, verified at that line by the operator, and
only the 8.21 s prediction and the "within 7%" comparison are the issue's alone.
The sentence now says that. 92712a0b0 is a merge of origin/main and carries no
change of its own.

…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]

…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 · 2026-08-21T17:34:58Z

`test_cpu_threadpool` is inherited, measured on `main` rather than assumed

main is red, so this pull request may only land if every failing job is
verified inherited against a main baseline. One case in build-test-cpu was
outside the set this repository already has evidence for, so it was chased rather
than waved through. Recording the measurement, because the next session to look at
a red build-test-cpu here should not have to redo it.

The case: tests/vt/test_cpu_threadpool.cpp:539, CHECK( ratio < 100.0 ), in
oversubscribed dispatch does not cost a scheduler timeslice.

It reds on main, six days before this branch existed. Scheduled baseline run
31896091307,
head_branch=main, event=schedule, at 5a0ffe9e372624a785af04683bdd8f9f41cf0693,
where it was the only failing test in the job (99% tests passed, 1 tests failed out of 484):

/home/runner/.../tests/vt/test_cpu_threadpool.cpp:536: MESSAGE: empty-op dispatch: 2 threads 0.421 us, 5 threads 47.139 us, ratio 111.969
/home/runner/.../tests/vt/test_cpu_threadpool.cpp:539: ERROR: CHECK( ratio < 100.0 ) is NOT correct!
  values: CHECK( 111.969 <  100 )

On this pull request it fails at the identical :499 / :536 / :539 sites, and the
decomposition is the same shape: the fitted arm is sub-microsecond (0.421 vs
0.371 us) while the oversubscribed arm is pinned at about 47 us (47.139 vs
47.592 us). That is a scheduler timeslice, not a code effect.

Three facts make it inherited rather than introduced:

It reds on main at a commit that predates this branch.
On both commits, the other two CPU jobs ran the same source and passed
the case (main: 1.56 s and 1.91 s; here: 1.43 s and 1.13 s). A verdict that
flips between sibling jobs on one commit cannot be a property of the diff.
This branch does not touch the test or the seam. The guard landed in
2028c8c74 (CPU backend: batch-1 decode is barrier-bound (47% sync), and paged attention branches per element #391) and its line numbers are byte-identical across both runs.

The build-test-cpu reds on main rotate. Over the 2026-08-12 to 2026-08-21
window, 59 main runs reached a failure verdict and 25 were fully green, so the
tree can sweep clean; the reds are drawn from a pool. By job count:
test_ltx2_video 11, test_nemotron_h_paged_forward 5, test_runner 2,
the test_minimax_music3_ar / test_ltx2_text_encoder / test_muse_glimmer_text
cluster 2 each, and test_cpu_threadpool 1. The sanitizer lanes add
test_resident_weight_host_addressable (9, tsan only) and test_ltx2_device (4),
and never failed the threadpool case in 49 job logs — it is confined to the fast
non-sanitized lane, which is consistent with a timing guard rather than a defect.

This is the same rotation #1052
describes when it notes that the identity of the failing test changes between runs
of an unchanged binary. It names test_cpu_threadpool as being on that list and
records that no issue named the test.

windows-msvc-cpu and windows-msvc-vulkan were checked the same way and are
unchanged: they fail at ERROR: CMakeLists.txt: MSVC /W4 /WX policy is negated on the C/C++ compile by /w, before any compilation, and open #1620
carries both failures today.

Bring the branch up to origin/main at 2e7f3be before the landing gate. The previous re-merge stopped at 369a0c2 and main has advanced six commits since, so the gate would otherwise run on a base the squash commit will not have. The only file both sides touch is .agents/issue-index.md. Its union driver merged the two appended regions, and the result was checked rather than trusted: `git diff --numstat` against origin/main reads `2 0`, so main's rows are preserved byte-for-byte and only this branch's two rows are added, and no issue id appears twice in the 536 rows. The four remaining paths on this branch are the muse-glimmer vision source, its test and two specs, none of which the delta touches. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

localai-bot · 2026-08-21T19:00:52Z

Landing blocked on #1629, not on anything in this change

This pull request is fresh-reviewed PASS, its one owed clause fix at .agents/specs/muse-glimmer-vision-attn-flash.md:151-152 is verified landed, it has been re-merged onto origin/main at 2e7f3bee7 (head 02be5f146), and test_cpu_threadpool — the one job the earlier session could not verify inherited — passed on the re-run, leaving test_runner.cpp:1557 as the whole of the build-test-cpu red, which is the known #1273 case owned by #1602 / #1608.

It is held anyway, because landing it turns main red on the required agent-record job once #1578 is in.

#1578 adds tests/scripts/test_check_attention_rung_consistency.py::ShippedTreeTests::test_the_population_is_not_empty, which asserts the scanned vt::Attention population is >= 9. The shipped tree has exactly 9 sites, so the floor has zero headroom, and this row deletes one of them. Measured on #1578's head with this branch's muse_glimmer_vision.cpp copied in: leaving the allowlist stem reds test_the_population_is_not_empty (8 not greater than or equal to 9), and deleting the stem reds that case and test_allowlist_holds_only_the_in_flight_stems. There is no green path for a removing row. The tree was restored byte-for-byte after each mutation and the suite returns to 31 tests ... OK.

The repair changes what a gate accepts, so under ## Changing the rules or a checker it needs its own row, spec and red-before evidence rather than an operator edit here. #1629 carries the detail and the two candidate directions.

Nothing in this branch needs to change. Re-verify the merge and land it once #1629 is repaired.

…the model tree, and the three rows it would have redded (#1629) `test_the_population_is_not_empty` was named for one guarantee and asserted another. The name promises non-emptiness; the assertion pinned `>= 9`, today's incidental number of `vt::Attention` call sites. The shipped tree has exactly 9, so the floor carried zero headroom and any row that legitimately REMOVED a naive-attention call turned the case red. Those rows are not hypothetical, and they are not strangers to this file: they are the three stems on scripts/attention-rung-allowlist.txt. The allowlist exists precisely so #1545 (muse_glimmer_vision) and the LTX-2.5 routing row (ltx2, ltx2_device) can reroute their calls without editing the lines they replace, and the checker backs that by reporting a cleaned-up stem as STALE rather than failing on it. The test then undid it. Composing this tree with #1579's muse_glimmer_vision.cpp, which routes that call to `vt::AttentionDenseFlash`, reds the case with `8 not greater than or equal to 9` - and the case runs in the required agent-record CI job, so `main` would have gone red on a change that did exactly what the allowlist invited. This is the shape AGENTS.md `## Records` names: never store a measurement of one file inside another file. A raw site total is a measurement of the model tree living in a test, and it couples every routing row to a line it does not own. The floor becomes `>= 1`, which is the guarantee the name always claimed and the one that actually matters - a scanner whose regex stops matching after a rename reports zero drift, and an empty scan and a clean tree file the same green. That is #1544's defect, and it stays covered. What the count was standing in for is covered without the coupling. The six deliberate sites are already pinned BY NAME, not by arithmetic, in `test_the_six_deliberate_sites_carry_a_marker`. The real risk a total never addressed is a bogus allowlist entry, so a new case asserts that every allowlisted stem names an existing model source under MODEL_DIRS. A typo is silent in both directions today: it excuses nothing, so the file it meant to cover goes on drifting unguarded, and the checker reports it only as STALE and exits 0. Verified - `muse_glimmer_vison` appended to the allowlist leaves `check-attention-rung-consistency.py` green at rc=0 and reds only the new case. The new case asserts FILE EXISTENCE, deliberately, and never scan membership. 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 the checker's own `stale_allowlist_entries` docstring states. Asserting the stem is still in `scanned` would rebuild the very lock this change removes. Nothing else moves. The checker's behaviour, the allowlist's stem set, and `test_shipped_tree_is_green`, `test_the_six_deliberate_sites_carry_a_marker` and `test_allowlist_holds_only_the_in_flight_stems` are untouched, so growth of the allowlist stays a review decision. Spec D7 still describes the tree accurately and needs no edit. Evidence, each mutation proven applied and restored by sha256: stubbing `scan_models` to `{}` reds the population case at `0 not greater than or equal to 1`; the typo'd stem reds the new case naming it while the checker stays green; and #1579's file composed over this tree runs 32/32 green with the checker at rc=0 printing `STALE (not a failure)` and the allowlist stem left in place, which is the point of the whole change. Preflight is green apart from `test_cpu_x86_llamacpp_floor`, the known load-dependent flake (#618); this box sat at load average 57. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

… the row that removes the LAST allowlisted stem (#1629) `test_the_ok_line_reports_the_excused_sites` did two jobs and stored a measurement of the model tree to do the first one. Its `assertGreater(excused, 0, "the shipped tree must exercise this branch")` was a live count of how many unmarked `vt::Attention` sites currently sit in files that scripts/attention-rung-allowlist.txt parks. That number is 3 today only because three rows are still in flight: #1545 (muse_glimmer_vision, PR #1579) and the LTX-2.5 routing row (ltx2, ltx2_device). Each of them exists to delete its own naive call, and the allowlist exists to let them do it without editing the lines they replace. When the LAST of the three lands the allowlist holds no stem, `excused` becomes 0, and this case reds while `check-attention-rung-consistency.py` is perfectly green at rc=0. It runs in the required agent-record CI job, so `main` would go red on a change that did exactly what the allowlist invited. That is the shape AGENTS.md `## Records` forbids -- never store a measurement of one file inside another file -- and it is the same shape as the `>= 9` population floor cab0b57 removed for this row, in the same not-yet-landed file. It is worth naming why it survived that repair: it does not fire for any of the three rows individually. With only #1579 landed the suite is 32/32 green. Only the third one trips it, which is the worst kind of lock to leave in a tree, because no in-flight change can find it. The two jobs are separated instead of weakened. The guard -- proving the checker's OK line actually EXERCISES the "unmarked and excused" branch with a non-zero count -- moves onto a tree this file constructs, so it holds forever regardless of what the model tree does. The mechanism is the smallest one that reaches `main()`: `mock.patch.object` over the two module-level names the report reads, `scan_models` and `ALLOWLIST`. The scan becomes a dict built by hand, exactly as every `MutationTests` case in this file already builds one, and the allowlist becomes a temporary file. A fixture directory of real .cpp sources was the alternative and is strictly more machinery for the same reach: `scan_models` computes `path.relative_to(ROOT)`, so a tempdir outside the repository raises, and a fixture dir inside it adds model sources to the tree the other cases scan. The constructed scan carries an unmarked site beside a marked one in the SAME allowlisted file, which is the case `sites - marked` cannot distinguish, and the assertion pins the whole OK line rather than a substring. A second constructed case pins the report at zero excused sites -- the state the allowlist exists to REACH. The checker prints the count even when it is zero, and its comment says so; nothing asserted it, and that is precisely the gap the floor was hiding. The shipped-tree job is kept as it was. `excused` is still RE-DERIVED from the tree and never pinned, so `assertIn(f"{excused} unmarked and excused by", ...)` holds at 3 today and at 0 after the last stem is cleaned up. Only the `assertGreater` line is gone. Nothing else moves: the checker, the allowlist, and every case this row's previous commit repaired are untouched. Evidence, each mutation proven applied by a diff and restored by sha256. RED BEFORE, simulating the end state on the unmodified tree -- the three naive calls routed to `vt::AttentionDenseFlash`, the allowlist emptied of stems, the pinned set in `test_allowlist_holds_only_the_in_flight_stems` set to `set()` as the landing row would -- the checker prints `0 unmarked and excused by 0 allowlisted in-flight stem(s)` at rc=0 while the case fails with `0 not greater than 0`. GREEN AFTER, that identical simulation with this change in place runs 34/34 with the checker still at rc=0. The guard still bites. Making the checker print `{0}` for `excused` reds the constructed case at `0 unmarked and excused by 1` against the expected `1`, and dropping the "unmarked and excused by" clause reds the zero case. The load- bearing run is the two composed: in the END STATE, with the checker's count broken, the shipped-tree case PASSES -- its recomputed `excused` is 0 and the broken line still says 0 -- and only the constructed case catches it. That is the coverage the deleted floor was standing in for, now held by something the model tree cannot switch off. Unmutated tree: 34 cases green, checker rc=0 printing `9 vt::Attention call site(s) in 9 model source file(s); 6 carry a recorded reason, 3 unmarked and excused by 3 allowlisted in-flight stem(s)`. Preflight is green apart from `test_cpu_x86_llamacpp_floor`, the known load-dependent flake (#618); this box sat at load average 44 on 20 cores. commit-trailers and commit-style SKIP because this base is behind origin/main, and were run directly over the range instead. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…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>

…h is #1578, and keep both sides' index appends GitHub reported this pull request CONFLICTING while `git merge-tree` against the same `origin/main` returned rc 0. The two disagree because `.agents/issue-index.md` carries `merge=union` in `.gitattributes`: local git honours the driver and unions the two appends, and the forge does not, so it sees two branches writing the file's last lines and calls it a conflict. #1578 appended four rows and this row appends two, which is exactly the shape that triggers it. Nothing else overlaps -- the merge reports `Auto-merging .agents/issue-index.md` and no other path. The union result is verified rather than trusted, because a union merge is precisely how a duplicated row gets in without a gate noticing. Against `origin/main` the file reads `2 0` in `git diff --numstat`: this row's #1545 and #1566 and nothing else, no deletions. No issue id appears twice in the 560 rows. `check-issue-index-append-only` is green on the committed merge. The rung interaction that held this row is settled and re-measured here. #1578 landed the #1629 repair, so `test_the_population_is_not_empty` asserts `>= 1` rather than the `>= 9` floor that a shipped tree of exactly nine sites left with no headroom. On the merged content `tests/scripts/test_check_attention_rung_consistency.py` runs 34 of 34 green and `scripts/check-attention-rung-consistency.py` exits 0, reporting `muse_glimmer_vision` as `STALE (not a failure)`. That stale entry is the designed outcome, not an oversight: `scripts/attention-rung-allowlist.txt` says the removing row may leave the deletion to whoever runs preflight next, and deleting the stem here would red `test_allowlist_holds_only_the_in_flight_stems` unless its expected set moved in the same change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…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>

mudler added 2 commits August 21, 2026 10:00

mudler added 4 commits August 21, 2026 10:36

mudler added 2 commits August 21, 2026 17:21

localai-bot mentioned this pull request Aug 21, 2026

check-attention-rung-consistency's population floor is pinned at today's exact count, so every allowlisted row reds it by removing its naive call #1629

Open

localai-bot mentioned this pull request Aug 21, 2026

fix(KERNEL-ATTN-DENSE-FLASH): a model on the naive attention kernel now says why, and AttentionDenseFlash advertises the head_dim it can launch (#1544) #1578

Merged

localai-bot merged commit 47a918d into main Aug 22, 2026
20 of 25 checks passed

This was referenced Aug 22, 2026

Three stems outlived their removing rows on the attention-rung allowlist, and a stale stem silently excused a DELETED marker in the file it named #1663

Closed

fix(#1663): a stale attention-rung stem excused a DELETED marker in the file it named #1665

Merged

Navigation Menu

fix(MUSE-GLIMMER-VISION-ATTN-FLASH): route the perception encoder off the correctness-grade attention kernel (#1545) - #1579

fix(MUSE-GLIMMER-VISION-ATTN-FLASH): route the perception encoder off the correctness-grade attention kernel (#1545)#1579
localai-bot merged 10 commits into
mainfrom
row/MUSE-GLIMMER-VISION-ATTN-FLASH

localai-bot commented Aug 21, 2026 •

edited

Loading

Uh oh!

localai-bot commented Aug 21, 2026

Uh oh!

localai-bot commented Aug 21, 2026

Uh oh!

localai-bot commented Aug 21, 2026

Uh oh!

localai-bot commented Aug 21, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

localai-bot commented Aug 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

How to verify

What is not verified

Uh oh!

localai-bot commented Aug 21, 2026

Review record, and the owed clause verified on the branch

The owed clause fix has already landed here

CI, read per job against a main baseline

Uh oh!

localai-bot commented Aug 21, 2026

Fresh review: PASS

Operator re-derivation

Uh oh!

localai-bot commented Aug 21, 2026

test_cpu_threadpool is inherited, measured on main rather than assumed

Uh oh!

localai-bot commented Aug 21, 2026

Landing blocked on #1629, not on anything in this change

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

localai-bot commented Aug 21, 2026 •

edited

Loading

CI, read per job against a `main` baseline

`test_cpu_threadpool` is inherited, measured on `main` rather than assumed