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)

[localai-bot]

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