Skip to content

fix(#1549): the LTX-2.5 DiT self-attention was on the correctness-grade kernel, and one forward cost 47.84 s - #1557

Merged
localai-bot merged 19 commits into
mainfrom
row/LTX25-DIT-ATTN-FLASH
Aug 22, 2026
Merged

fix(#1549): the LTX-2.5 DiT self-attention was on the correctness-grade kernel, and one forward cost 47.84 s#1557
localai-bot merged 19 commits into
mainfrom
row/LTX25-DIT-ATTN-FLASH

Conversation

@localai-bot

@localai-bot localai-bot commented Aug 21, 2026 •
edited
Loading

Copy link
Copy Markdown
Collaborator

One LTX-2.5 DiT forward at 768x448/49f measured 47.84 s on GB10 -- n=119, median 47.91 s, spread 5.8%, from the engine's own last= samples -- and first-order arithmetic says it should take well under a second.

The DiT self-attention called vt::Attention, which on CUDA resolves to the kernel whose own header calls itself "Correctness-grade (M0.9)": one 256-thread block per (query, head), a 256-wide shared-memory tree reduction per key, and no K/V tiling, so K and V are re-read from global once per (query, head). At that geometry the video stream is 2352 tokens x 32 heads = 75,264 blocks each looping 2352 keys, over 48 layers.

The attribution is arithmetic rather than assertion. .agents/specs/multimodal-speed.md measures that same kernel on that same box at 5.70 ns per block-key iteration, and 1.77e8 x 5.70 ns x 48 = 48.4 s against the measured 47.84 s. A 1% match, which leaves the other 1% for 48 blocks of GEMMs, norms, RoPE, gating and six attentions each.

After this change one forward is 7.680 s (n=19, median, 768x448/49f, GB10). The ratio is ~6.0x, and it is not an A/B. Both qualifications are load-bearing and both are under "The A/B" below.

Why nobody saw it, which is the part worth keeping

kAttention is deliberately frozen on the naive kernel so that text decode stays byte-identical. That decision is correct and this change does not touch it.

The consequence is the defect. The fast kernels are separate ops that each caller must opt into by name -- kAttentionDenseFast, kAttentionDenseFlash, kAttentionDenseFa2. There is no automatic selection, no shape routing, and no fallback notice. A model that never opts in gets correct output at roughly 500x the cost, with no warning anywhere.

Nothing in this tree can detect that. The output is right, so every golden passes. The op is registered, so no refusal fires. GetOpProviderStats counts the naive selection and reports it as a success, because it is one. The only symptom is a wall clock, and a diffusion render has no reference wall clock to be held to. That is why the gate for this change is a dispatch observation and not a number.

What changed

ltx2_device.cpp calls vt::AttentionDenseFlash in the self-attention branch. That is the whole of the product change. The dispatch RULE is unchanged: the branch is still chosen by context == nullptr && a.bias == nullptr, upstream's own self-attention marker, never by what the numbers happen to be. Only the op it calls moves. Its square-problem contract holds by construction, because that branch is entered only when s == tq.

VLLM_LTX2_DIT_FLASH_ATTN=0 restores the old op, so both arms of the measurement run from one binary. Same shape as VT_FA2_DENSE, documented in docs/ENVIRONMENT.md as a measurement lane and never a configuration. Both remaining naive vt::Attention call sites -- that one, and ltx2.cpp's CPU-only host arm -- carry a // VT-ATTN-NAIVE: line saying why, the form #1578's checker defines for a deliberate site.

The host arm at ltx2.cpp is deliberately NOT moved. It computes into std::vector<float> and is CPU-only by construction, and on CPU both ops are the same registered function (src/vt/cpu/cpu_ops.cpp:3750-3761), so the swap would be a byte-identical no-op that moves the L2 parity reference off the reference op.

The shared-memory cap-raise is REVERTED, and that is the change since the review

An earlier version of this branch also raised LaunchAttentionDenseFlash's dynamic shared-memory cap through cudaFuncSetAttribute, moved SetDynamicSmemOptIn onto a shared seam, and added head_dim contract cases to tests/vt/test_ops_attention.cpp. All of it is gone. src/vt/cuda/cuda_ops.cu, cuda_device_caps.h, cuda_arch_tactics.cu, cuda_paged_attn.cu and tests/vt/test_ops_attention.cpp are byte-identical to main.

The swap never needed it. LTX renders at the stream dtype, and in production that is bf16 (ltx2_device.cpp:1183). The flash op's K/V tile is 2 * kFlashBc(64) * head_dim * sizeof(Tin), so the video stream's head_dim 128 asks 32,768 B and the audio stream's head_dim 64 asks 16,384 B. Both are inside the 49,152 B every CUDA architecture gives a launch without any opt-in at all.

What the raise was actually serving does not fit. The shapes over 48 KiB here are f32: head_dim 128 at 65,536 B, and head_dim 256 at 131,072 B against GB10's queried ceiling of 101,376 B (cuda_device_caps.h:46). The 256 shape therefore does not fit even with the opt-in -- it was falling back to the bit-identical AttentionDenseFast while the case reported a launch. So the raise bought this row one shape it does not run, and paid for it by moving a shared helper across two files and colliding with #1578 on the same lines.

#1578 owns the bound and takes the opposite, better approach. Rather than raising the cap it makes the ADVERTISED domain honest: AttentionDenseFlash declares head_dim <= 256 while it can only launch bf16 192 / f32 96, and #1578 narrows the declaration to what the code can do and refuses above it. That is a property of the code rather than of whichever device is underneath, and it is the supports_head_size() polarity vLLM already has. #1578 merges first, and after it bf16 head_dim 128 is inside the declared bound, so this change is unaffected. There is now nothing left to conflict.

One consequence is disclosed rather than left to be found. With the raise gone, the f32 L2 parity arm at production geometry (head_dim 128, 65,536 B) reaches AttentionDenseFlash and cannot launch: a cudaGetLastError throw at cuda_ops.cu:3352 today, a VT_CHECK naming the head_dim once #1578 lands. It fails loud in both worlds and never silently, and nothing gated reaches it -- production is bf16, and the f32 arm is a parity reference exercised at the fixture's reduced dimensions. Filed under ## Owed.

Numerics, measured rather than asserted

On CPU: byte-identical. kAttention and kAttentionDenseFlash are the same registered function pointer, and the goldens are unmoved -- test_ltx2_device 22/22, 652/652; test_ltx2 43/43, 4581/4581; test_ltx2_video 102/102, 4194/4194; test_ops_attention_cross 9/9, 32/32.

On CUDA: NOT bit-identical, and here is the number. The warp kernel groups the head_dim partial sums across 32 lanes instead of a 256-thread block, so the same f32 online softmax associates differently. On dgx:gpu0 the host-vs-device parity case measures video 8.9407e-08, audio 4.47035e-08 against its committed 2e-5, and bf16 CUDA-vs-CPU-backend at 0. That is f32 round-off scale, 224x inside the gate, and the gate was not widened.

That is the only numeric evidence there is, and it bounds less than it looks like it bounds. The case runs the fixture's reduced dimensions, so it bounds the ARITHMETIC change -- a length-D sum reassociated -- and not the change at head_dim 128 with 2352 keys over 48 layers. A diffusion render has no token gate to fall back on, and the flash arm was interrupted before writing any frames, so no pixel comparison exists either, not even against the completed 49-frame baseline render already on the NAS. Filed as #1612 and listed under ## Owed.

The test_ops_attention evidence is WITHDRAWN, not restated. The GB10 lease ran it at 10/10 and 88,439 assertions, and this change claims nothing from that run: it measured the head_dim cases that came with the cap-raise, and those are no longer in the tree. Two of its three arms would not have supported the claim anyway, for the reason above. The test_ltx2_device rows and the render below DO survive the revert, and that is arithmetic rather than assertion -- the measured binary carried the opt-in call, but the helper returns immediately below 49,152 B, so it was a no-op on every launch those numbers came from.

Reachability, twice

Unit. A new case drives the production entry point Ltx2DitForwardDevice -- called from the denoise loop at ltx2_video.cpp:4246 -- and asserts the dispatch two-sidedly through GetOpProviderStats: kAttentionDenseFlash selected exactly 8 times (two self-attentions x two blocks x two batch rows) and kAttention selected 0. The negative half is what makes it a routing proof rather than an addition proof. The case scope-guards its own process state, because it enables a counting instrument and sets an env var and contains REQUIREs; measured with a scratch REQUIRE(false) and an appended observer, without the guards the observer reads the knob still set and 8 leaked selections, and with them it reads neither.

That :4246 is itself a repair. This branch and its spec both cited ltx2_video.cpp:4055-4059, which points at prose about the res_2s step counter -- at the row's declared base 6b48edb2c and at HEAD alike, so re-reading it at either revision would have caught it. It is the citation the whole reachability argument rests on. cpu_ops.cpp:3551-3562 was wrong the same way and pointed at FusedStore. Every other anchor in the spec has been re-read against the base rather than carried forward, and eight more were corrected.

Mutation M1: restore vt::Attention at the call site, +1/-1, compile rc 0. Both halves went red -- CHECK( 0 == 8 ) and CHECK( 8 == 0 ), exit 1 -- while every golden case in the same binary stayed green. That contrast is the finding: no numerical gate in this tree can see a 500x slower kernel that computes the right answer. Tree restored and re-gated green.

Production, on the real model. With VT_OP_PROVIDER_STATS=1 the full 21.00B render at 768x448/49f on GB10 announces op=21 device=1 (kAttentionDenseFlash on CUDA) and announces op=18 device=1 (kAttention on CUDA) zero times. Same two-sided claim, taken through --device cuda at full scale rather than on a fixture.

The A/B: one arm measured, the pair still PENDING

Lease 6c724dfd on dgx:gpu0, source 30dce3a1d, one binary built in-lease with cutlass-nvfp4, cutlass-fp8 and FA-2 all ENABLED for [121a]. Correctness cleared before any speed number was read.

Flash arm, per DiT forward at 768x448/49f = 2352 tokens, from the engine's own last= lines:

n median mean min max spread
19 7.680 s 7.633 s 7.109 s 8.196 s 14.2%

The naive arm did not run, so this is not an A/B. At forward 20 the rc worker was lost and dgx:gpu0 read unhealthy (no contact). The cause is UNPROVEN and this change does not name one: no memory trace was taken and the box did not return to be asked. What is established is that the harness as first written carried no memory guard and no sample cap, which is a defect in this row's harness rather than a finding about the change.

The 47.84 s denominator carries two confounds this arm does not, and both inflate the ratio. Neither was disclosed before:

  • A stack sampler. The denominator ran under runguard.py --stack-period 12 (render.log:1), which eu-stacks the process and so ptrace-stops every thread. Its own stacks.txt prices that: 523 samples, median inter-sample delta 12.40 s against a 12.0 s period, so ~0.40 s median and 1.50 s max of stopped process per sample. About 3.9 samples land inside each 47.84 s forward, ~1.54 s, ~3.2%. Correcting only the denominator gives 46.3 s / 7.680 s = 6.03x.
  • A different prompt. render.log:1 carries a ~70-word prompt; the harness uses one short sentence, and ltx2_video.cpp:2253 sets context_tokens = encoded.seq unpadded, so the DiT's cross-attentions see a different number of keys in each arm. Corroborated rather than inferred: conditioning.tower is 45.013 s against 28.426 s. Same sign, and not quantified.

So the defensible statement is the range 6.03x to 6.23x, quoted as ~6.0x, with the sampler correction named and the prompt confound uncorrected and pushing the same way. 6.23x survives in the records only as the uncorrected upper end of that range, never on its own. The A/B gate reads PENDING and this change does not claim otherwise.

The flash arm's artifacts do not record what it ran, which is why those confounds had to be established from a phase duration. arm-flash.log opens at [render] + load with no command line, wd-flash/ is empty, no phase-log.json was written, and the only description of the run was a mutable NAS path edited 25 minutes after it finished. Both halves are repaired: the harness is committed as scripts/ltx25-dit-attn-flash-ab.sh, and every arm now writes its own invocation -- harness sha256, binary sha256, source SHA, geometry, seed, prompt, resolved command line -- to line 1 of its own log. It also caps each arm at 13 samples, holds a 12 GiB MemAvailable floor, caches the build on the source SHA, and runs the naive arm first.

That harness also had a precondition that could never pass. It grepped cuda_ops.cu for FlashTileSmemOptIn, a spelling no revision of this change ever used, so it counted 0 and would have exit 42-ed on a correct tree as readily as on a wrong one. It is removed with the cap-raise it guarded.

CI

Owed, filed and not folded in

All three are listed under ## Owed in .agents/specs/ltx25-dit-attn-flash.md, together with the two scripts/attention-rung-allowlist.txt stems that #1578's checker will report STALE once the markers here meet it.

Closes #1549

FOLLOWING_AGENTS_PROTOCOL

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

@localai-bot localai-bot mentioned this pull request Aug 21, 2026
Closed
mudler added 11 commits August 21, 2026 15:31
@mudler
spec(LTX25-DIT-ATTN-FLASH): the LTX-2.5 DiT never opted into a fast a…
899450a 
…ttention op, and paid 47.84 s a forward (#1549)

One LTX-2.5 DiT forward at 768x448/49f measures 47.84 s on GB10 (n=119, median
47.91 s, spread 5.8%), and first-order arithmetic says it should take well under
a second. The cause is that `ltx2_device.cpp:421` calls `vt::Attention`, which on
CUDA is the kernel whose own header at `cuda_ops.cu:1456-1459` calls itself
"Correctness-grade (M0.9)": one 256-thread block per (query, head), a 256-wide
shared-memory tree reduction per key, and no K/V tiling.

The attribution is arithmetic rather than assertion. `multimodal-speed.md:24-26`
measures that same kernel on that same box at 5.70 ns per block-key iteration,
and LTX's 1.77e8 block-key iterations per call over 48 layers come to 48.4 s
against the measured 47.84 s.

The reusable half is why nobody saw it. `kAttention` is deliberately frozen on
the naive kernel so text decode stays byte-identical, and the fast kernels are
separate ops each caller must name. There is no shape routing and no fallback
notice, so a model that never opts in gets correct output at roughly 500x the
cost with no warning anywhere: the goldens pass, no refusal fires, and the
provider stats count the naive selection as the success it is.

This commit is the spec only, so the implementation that follows it can be read
against a statement of intent that predates it. It also files two things found
while writing it rather than folding them in: FA-2 refuses head_dim 128 (#1551),
and the same defect shape reaches every other `vt::Attention` caller (#1552).

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
fix(#1549): route the LTX-2.5 DiT self-attention onto the flash-tiled…
31b085c 
… op, and give the tile the shared memory it asks for

The DiT self-attention called `vt::Attention`, which on CUDA resolves to the
kernel whose own header calls itself "Correctness-grade (M0.9)": one 256-thread
block per (query, head), a 256-wide shared-memory tree reduction per key, and no
K/V tiling. At 768x448/49f that is 75,264 blocks each looping 2352 keys, over 48
layers, and one DiT forward measured 47.84 s on GB10.

Three changes, and the third is a prerequisite rather than a bonus.

`ltx2_device.cpp` calls `vt::AttentionDenseFlash` in the self-attention branch.
The dispatch RULE is unchanged: the branch is still chosen by
`context == nullptr && a.bias == nullptr`, upstream's own self-attention marker,
and never by what the numbers happen to be. Only the op it calls moves.
`VLLM_LTX2_DIT_FLASH_ATTN=0` restores the old op so the A/B runs on one binary,
the shape `VT_FA2_DENSE` already takes.

`cuda_ops.cu` opts the flash K/V tile into more than the 48 KiB of dynamic shared
memory every CUDA architecture guarantees without asking. The op advertises
head_dim up to 256; at f32 that tile is 64 KiB at head_dim 128 and 128 KiB at
256, so those two shapes had never launched at all, and the symptom was a bare
`invalid argument` from a kernel that had not run. The ceiling is QUERIED through
the same cached seam `cuda_paged_attn.cu` uses, not assumed, and a tile that
cannot be opted into falls back to `AttentionDenseFast` rather than refusing --
the flash kernel's own contract is that the two are bit-identical, so the
fallback costs the tiling and nothing else. LTX's f32 arm is a supported arm at
head_dim 128, so without this the swap would have turned slow-but-correct into a
refusal.

The host arm at `ltx2.cpp` is deliberately NOT moved. It computes into
`std::vector<float>` and is CPU-only by construction, and on CPU both ops are the
same registered function, so the swap would be a byte-identical no-op that moves
the L2 parity reference off the reference op.

NUMERICS, stated rather than glossed. On CPU this is byte-identical and the
goldens are unmoved: `test_ltx2_device` 22/22 and 652/652, `test_ltx2` 43/43 and
4581/4581. On CUDA it is NOT bit-identical -- the warp kernel groups the head_dim
partial sums across 32 lanes instead of a 256-thread block, so the same f32
online softmax associates differently. The binding gate is the existing
host-vs-device parity case, at its committed 2e-5 f32 and 5e-3 bf16, and that
case is not widened by this change.

REACHABILITY. A new case drives the production entry point
`Ltx2DitForwardDevice` and asserts the dispatch two-sidedly through
`GetOpProviderStats`: `kAttentionDenseFlash` selected exactly 8 times (two
self-attentions x two blocks x two batch rows) and `kAttention` selected zero.
Mutation M1 restored `vt::Attention` at the call site (+1/-1, compile rc 0) and
both halves went red -- `CHECK( 0 == 8 )` and `CHECK( 8 == 0 )`, exit 1 -- while
every golden case in the same binary stayed green, which is the whole point: no
numerical gate in this tree can see a 500x slower kernel that computes the right
answer.

The two things found while doing this are filed and NOT folded in: FA-2 still
refuses head_dim 128 (#1551), and the same opt-in-by-name defect reaches every
other `vt::Attention` caller (#1552). Both are listed under `## Owed` in the
spec.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
spec(LTX25-DIT-ATTN-FLASH): the shared-memory failure this row fixes …
4979228 
…is already documented one file over (#1549)

`src/vt/cuda/cuda_attention_cross.cu:88-96` states the same defect, with the same
arithmetic, for the same model: at f32 and head_dim 128 a 64-column K/V tile asks
for 64 KiB of dynamic shared memory, over the 48 KiB a launch gets without opting
in, "so the kernel would fail to launch on exactly the real geometry while every
reduced-dimension gate (head_dim 8 and 4) passed". That kernel is itself a port of
`AttentionDenseFlashKernel`; the problem was solved in the copy and never carried
back to the original. This turns section 4.3 from a reading of the code into a
corroborated finding, and it is why the new `test_ops_attention` case runs at the
real head_dims rather than the fixture's.

It also tightens section 2's attribution rather than loosening it: because
`vt::AttentionCross` on CUDA is already flash-tiled, the DiT's four
cross-attentions per block were never on the naive kernel, and only the two
self-attentions were slow.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
test(#1549): flip the A/B knob through the shared shim, not a private…
b4f00c3 
… _WIN32 branch

The reachability case set `VLLM_LTX2_DIT_FLASH_ATTN` with its own
`#ifdef _WIN32 / _putenv_s / setenv` block. That is the exact thing issue #603
exists to stop: `setenv` and `unsetenv` are POSIX, MSVC has neither, and three
files had already grown a private copy before `tests/support/test_env.h` landed
to hold it once. Nothing would have caught it either, because
`check-windows-portability.py` scans no test translation unit at all (#1107), so
the first report would have been the `windows-msvc-*` lane going red for a fourth
reason.

`vllm_test::SetEnv` and `UnsetEnv` also throw rather than return a status, which
is the behaviour this case wants: a test whose environment did not take effect is
not running the arm it claims, and here the arm IS the assertion.

Focused gate unchanged and green: 22/22 cases, 652/652 assertions, with the
reachability case still reading `kAttentionDenseFlash selections = 8 (want 8),
kAttention selections = 0`.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
record(LTX25-DIT-ATTN-FLASH): the GB10 gates passed and the flash arm…
3463a50 
… is 7.680 s, but the worker died before the naive arm (#1549)

Lease `6c724dfd` on `dgx:gpu0`, source `30dce3a1d`, one binary built in-lease with
cutlass-nvfp4, cutlass-fp8 and FA-2 all enabled for `[121a]`.

CORRECTNESS FIRST, and it passed before any speed number was read.
`test_ops_attention` 10/10 with **88,439** assertions against 23 on a CPU box, so
the new f32 head_dim 64/128/256 case really ran and the shared-memory launch
failure is repaired against a device rather than against a reading of the code.
`test_ltx2_device` 22/22 with 749 assertions. The numerics claim is now a number:
the swap is NOT bit-identical on CUDA and the measured deviation is
**8.94e-08 video / 4.47e-08 audio**, f32 round-off scale, 224x inside the 2e-5 the
case was already held to. The gate was not widened.

REACHABILITY AT FULL SCALE, from the render's own log rather than a fixture.
With `VT_OP_PROVIDER_STATS=1` the 21.00B render at 768x448/49f announces
`op=21 device=1` (`kAttentionDenseFlash` on CUDA) and announces `op=18 device=1`
(`kAttention` on CUDA) **zero** times. Same two-sided claim the unit case makes,
taken through the production entry point on the real model.

THE FLASH ARM IS MEASURED: n=19, median **7.680 s**, mean 7.633, 7.109 to 8.196,
spread 14.2%, reduced from the engine's own `last=` lines. Against the recorded
47.84 s that is 6.23x, inside the predicted range.

THE NAIVE ARM DID NOT RUN, so there is no same-binary A/B. At forward 20 the `rc`
worker was lost and `dgx:gpu0` read `unhealthy (no contact)`. GB10 shares host RAM
with the GPU and an unconstrained job has OOM-rebooted this box before, and
`ab.sh` as first written carried no memory guard and no sample cap. That is a
defect in this row's harness, not a finding about the change, and the A/B gate
therefore reads PENDING rather than satisfied: 6.23x is a cross-run comparison
against a number a different binary produced in a different lease, which is the
weaker form the same-binary rule exists to replace.

The harness now caps each arm at 13 samples, holds a 12 GiB `MemAvailable` floor,
traces memory per arm, caches the build keyed on the source SHA so a resumed run
does not re-spend 18 minutes compiling, and runs the NAIVE arm first. The
previous order took the cheap arm first and lost the box before the expensive
one, which is how a two-arm measurement became a one-arm one.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
record(LTX25-DIT-ATTN-FLASH): score the measured arm and the PENDING …
2930a6b 
…one, on the public pages this claim reaches (#1549)

The GB10 run produced one arm and no pair, so the record has to say both
things: 7.680 s per DiT forward at `768x448/49f` (n=19, from the engine's own
`last=` lines) is measured, and the A/B against the 47.84 s naive denominator
is `PENDING` because the worker was lost at forward 20 and the naive arm never
ran. The cause of that loss is UNPROVEN and this record refuses to name one:
no memory trace was taken and the box did not return to be asked.

`docs/STATUS.md` and `docs/BENCHMARKS.md` are written in the same commit as
`.agents/benchmark-record.md`. A new measurement is a claim about the project,
and `scripts/check-doc-checkpoint.py` holds every commit that records one to
the surfaces that claim reaches a reader through. The two record commits this
one replaces wrote the record without the surfaces, which is what put
`documentation-checkpoint` red on this branch.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
fix(#1549): the flash K/V tile asks the device for its bound, and a s…
bd0f58f 
…hape that does not fit is refused by name

Review finding on this branch: the first version of this fix answered a tile
that does not fit by falling back to `AttentionDenseFastKernelCuda`, on the
argument that the flash kernel's own header makes the two bit-identical by
construction, so the fallback is numerically free. Both halves of that are
true and it is still wrong. A silent degradation to the untiled rung is the
DEFECT THIS ROW EXISTS TO REMOVE -- #1549 is entirely about a caller sitting on
a 500x slower kernel with nothing anywhere saying so -- and because the two
kernels agree bit-for-bit, no numeric gate anywhere can tell a fallback from a
launch. That is not hypothetical: `test_ops_attention` reported 10/10 and
88,439 assertions on GB10 while flash never launched at f32 head_dim 256 once.

So the launcher refuses instead, naming the device and the shortfall.

The BOUND is derived at runtime and is written down nowhere. Against GB10's
queried 101,376 B, with `kFlashBc = 64` and a tile of `2 * 64 * d * sizeof(T)`:

                no opt-in (49,152 B)    opt-in, GB10 (101,376 B)
      bf16      d <= 192                d <= 256
      f32       d <=  96                d <= 192

Neither column belongs in the code. 192/96 caps GB10 at two thirds of the
head_dim it can serve; 256 assumes an opt-in ceiling GB10 does not have, and
f32 at 256 wants 131,072 B. `cuda_device_caps.h:46` already says the rule --
"other architectures differ and MUST be asked, not assumed" -- so the tile is
asked about per device. The `d <= 256` VT_CHECK stays: that one is the kernel's
own `kMaxPerLane = 8` register bound and is a property of the code, not of any
device.

`SetDynamicSmemOptIn` is not a new function and not a second copy. It was
cuda_paged_attn.cu's file-local helper; it moves to the shared seam it always
belonged to -- declared in `cuda_device_caps.h`, defined beside
`DynamicSmemFits` in `cuda_arch_tactics.cu` -- because a second copy is the
parallel path AGENTS.md "Shared seams" forbids, and the two would have diverged
immediately, the second one having been written to fall back silently. Its six
existing call sites are unchanged and GB10 takes the same branch as before.

The test is retitled to what it exercises. f32 head_dim 128 (65,536 B) is the
shape that was red and is repaired; f32 head_dim 256 is NOT repaired and cannot
be on this device. Its arm now asserts the device-derived outcome -- launch and
match where the ceiling reaches 131,072 B, refuse and name BOTH numbers where it
does not -- rather than a numeric match that a bit-identical fallback satisfies
either way.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
test(#1549): the two naive attention call sites say why they stay, an…
4afe3fa 
…d the reachability case restores its own instrument

Three review findings, all in the same two files plus their test.

TWO ANCHORS WERE WRONG, and one of them is load-bearing. The reachability
case's ENTRY POINT paragraph and the spec's §6 both cited
`src/vllm/multimodal/ltx2_video.cpp:4055-4059` as the production call to
`Ltx2DitForwardDevice`. At HEAD and at this row's base that region is prose
about the res_2s step counter; the call is at `:4246`. That is the citation the
whole reachability argument rests on. The second, `cpu_ops.cpp:3551-3562`, is
`FusedLoad`/`FusedStore`; the registrations it meant are at `:3750-3761`. The
claim there was true and only the anchor was wrong.

THE CASE LEAKED PROCESS STATE ON ANY UNWIND. It enables
`vt::EnableOpProviderCallStats` and sets `VLLM_LTX2_DIT_FLASH_ATTN`, both
per-process, and it contains `REQUIRE`s. A failed `REQUIRE` or a throw from
either forward left the counting instrument enabled and the knob set for the
other 21 cases in the binary, so an unrelated case's result would depend on
which case failed first. Two scope guards now hold both. Measured rather than
asserted: with a scratch `REQUIRE(false)` after the `SetEnv` and an observer
case appended after it, the observer without the guards reads the knob still
set and counts 8 leaked `kAttention` selections; with them it reads neither.
Tree restored byte-for-byte and re-gated 22/22, 652/652.

BOTH NAIVE `vt::Attention` CALL SITES NOW RECORD THEIR REASON IN THE FILE, in
the `// VT-ATTN-NAIVE:` form #1578's allowlist defines for a deliberate one.
This change ADDS one (the A/B knob's off arm in `ltx2_device.cpp`) and leaves
`ltx2.cpp`'s host arm in place, while that allowlist exempts both stems on the
ground that "another row is removing" their naive call. This row does not
remove either, so without the markers both files would be permanently and
silently exempt from the checker that row lands.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
bench(#1549): commit the A/B harness, and make every arm state its ow…
1793a8e 
…n invocation

The flash arm's 7.680 s is not reproducible from its own evidence.
`arm-flash.log` opens at `[render] + load` with no command line, `wd-flash/` is
empty and no `phase-log.json` was written, so the geometry, prompt, seed and
sample cap behind that number cannot be read off the run. The only description
of it was `/mnt/nas_share/rc/ltx25-attnflash/job/ab.sh`, a mutable path on a
share, whose mtime is 25 minutes AFTER the run finished and which is not
committed anywhere. The 47.84 s denominator has its full command line as line 1
of its own `render.log`; that asymmetry is why this branch's two measurement
confounds had to be established from a `conditioning.tower` duration instead of
being read off a recipe.

Two repairs. The harness is committed here, so a revision of it is immutable and
citable. And each arm now writes its own invocation to line 1 of its own log --
harness sha256, binary sha256, source SHA, geometry, seed, prompt and the
resolved command line -- with a `harness_sha256` line into `PROVENANCE` beside
the binary's. `set -x` is deliberately not used: it would interleave the trace
with the engine's progress lines, which are the samples being reduced.

Committed as-is otherwise. The file already carries the sample cap, the
`MemAvailable` floor, the per-arm memory trace, the SHA-keyed build cache and
the naive-arm-first ordering that the lost-worker run lacked, and rewriting it
in the same change would leave nothing to compare the next run against.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
record(LTX25-DIT-ATTN-FLASH): the 256 launch never happened, the 6.23…
25d294e 
…x carried two confounds, and neither was disclosed (#1549)

Review findings against the records, all of them narrowing a claim.

f32 HEAD_DIM 256 WAS NEVER REPAIRED, and three surfaces said it was. That tile
is `2 * 64 * 256 * 4` = 131,072 B against GB10's queried 101,376 B, so it never
fitted; the launcher fell back to `AttentionDenseFast`, which is BIT-IDENTICAL
to the flash kernel by contract, and all 20,480 of that arm's assertions passed
with flash never launching at 256 once. A numeric comparison cannot separate the
two -- that IS the contract -- so the case measured a fallback and reported it
as a launch. head_dim 128 (65,536 B) is genuinely repaired and that claim stands.

THE RATIO IS ~6.0x, NOT 6.23x, and the correction is named rather than quietly
applied. Two confounds, both inflating it, neither disclosed:

  * The denominator ran under `runguard.py --stack-period 12`, which `eu-stack`s
    the process and so ptrace-stops every thread. Its own `stacks.txt` prices
    that: 523 samples, median inter-sample delta 12.40 s against a 12.0 s
    period, so ~0.40 s median and 1.50 s max of stopped process per sample.
    ~3.9 samples land inside each 47.84 s forward, ~1.54 s, ~3.2%. Correcting
    only the denominator: 46.3 s / 7.680 s = 6.03x. The flash arm had no sampler.
  * The two arms used different prompts. The denominator's `render.log:1` has a
    ~70-word prompt and the harness has one short sentence, and
    `ltx2_video.cpp:2253` sets `context_tokens = encoded.seq` UNPADDED, so the
    DiT's cross-attentions see a different number of keys. Corroborated:
    `conditioning.tower` 45.013 s against 28.426 s. Same sign, not quantified.

The defensible statement is the range 6.03-6.23x with the sampler correction
named and the prompt confound uncorrected and pushing the same way. `~6.0x` is
what the pages now quote.

WHAT ELSE IS RECORDED: the runtime-derived bound and why neither #1578's 192/96
nor this branch's earlier 256 is it, with the note that #1578 must reconcile
onto it and that this row does not edit that branch; the two corrected anchors;
the scope-guard mutation and its counter-proof; the flash arm's missing
invocation and the committed harness that replaces it; and the
`documentation-checkpoint` red, which was this branch's own and not inherited,
together with the side effect that its `set -eu` step stopped
`check-now-current.py` and `check-role-discipline.py` from running in CI at all.

NEWLY OWED, filed as [#1612](#1612):
there is no numeric or pixel comparison at production geometry. The only numeric
gate is the reduced-dimension host-vs-device case, which bounds the arithmetic
change and not the change at head_dim 128 with 2352 keys over 48 layers; a
diffusion render has no token gate; and the flash arm was interrupted before
writing any frames, so no pixel A/B exists even against the completed 49-frame
baseline render already on the NAS.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
merge(row/LTX25-DIT-ATTN-FLASH): take origin/main 9e4e04a under the b…
b97910e 
…ranch

The branch is rebuilt on e2a9e03 and main has moved again since. Merging
rather than rebasing keeps the repair commits reviewable against the head the
fresh review read.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@localai-bot
localai-bot force-pushed the row/LTX25-DIT-ATTN-FLASH branch from 2f39a94 to b97910e Compare August 21, 2026 16:03
@mudler
docs(#1549): the paged-attention tile comment no longer says "above" …
2b80b4c 
…of a function that moved

`SetDynamicSmemOptIn` left this translation unit for the shared seam in the
preceding commit, so `cuda_paged_attn.cu:944`'s "enforced by SetDynamicSmemOptIn
above" points at nothing. The symbol name is what survives the move, and the
note at :96 says where it went.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot pushed a commit that referenced this pull request Aug 21, 2026
@mudler
record(KERNEL-ATTN-DENSE-FLASH): D2's rejection of the cap raise is a…
2859fc7 
… measurement, not a convenience (#1544)

D2 rejected `cudaFuncSetAttribute` partly because "it cannot be verified without
a device, and this row has no lease". That reads as a preference. The number is
now available from #1557's review: GB10's queried opt-in ceiling is 101,376
bytes, and head_dim 256 in f32 wants 131,072. The raise therefore cannot make the
advertised 256 true for f32 on the part this project gates on, which is the exact
width that motivated it, so narrowing beats raising on measured grounds.

#1573 stays owed and says why the new number does not discharge it: 101,376 is a
device value that bounds what an opt-in could buy, and it proves nothing about
whether the launcher's refusal executes. That still needs a lease.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@localai-bot localai-bot mentioned this pull request Aug 21, 2026
Merged
mudler added 6 commits August 21, 2026 18:31
@mudler
merge(row/LTX25-DIT-ATTN-FLASH): take origin/main c477f6f under the b…
6b0ec47 
…ranch

`.agents/benchmark-record.md` was the only conflict and both sides are pure
appends: main's `### The Qwen3.5-4B 1.0283x row ran against an UNFUSED
denominator` subsection and this row's `## LTX25-DIT-ATTN-FLASH` top-level
section. Resolved by taking main's version byte-for-byte first -- it is a `###`
under the section above it and would be orphaned by a `##` inserted before it --
then this row's section. Neither side lost a line.

`.agents/issue-index.md` union-merged clean at `4 0`: four appended rows, no
deletion, no duplicated id.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
revert(#1549): drop the shared-memory cap-raise, which this row never…
4620230 
… needed

LTX-2.5's DiT renders at the stream dtype, and in production that is bf16
(`ltx2_device.cpp:1166`). At head_dim 128 the flash op's K/V tile is
`2 * kFlashBc(64) * 128 * sizeof(bf16)` = 32,768 B, and the audio stream's
head_dim 64 asks 16,384 B. Both are inside the 49,152 B of dynamic shared memory
every CUDA architecture gives a launch WITHOUT an opt-in, so the routing this row
exists for never needed `cudaFuncSetAttribute` at all.

What the raise was actually serving was f32: head_dim 128 at 65,536 B, and
head_dim 256 at 131,072 B against GB10's queried ceiling of 101,376 B
(`cuda_device_caps.h:46`) -- which does not fit even WITH the opt-in, and was
being answered by the bit-identical `AttentionDenseFast` while the case reported
a launch. So the raise bought one shape this row does not run, and paid for it by
moving a shared helper across two files and colliding with #1578 on the same
lines of `cuda_ops.cu` and `test_ops_attention.cpp`.

#1578 is the correct treatment and it merges first: instead of raising the cap it
narrows the ADVERTISED head_dim domain to what the code can launch and refuses
above it, which is a property of the code rather than of whichever device is
under it. After it, bf16 head_dim 128 is inside the declared bound and this row's
swap is untouched.

Reverted in full. `src/vt/cuda/cuda_ops.cu`, `cuda_device_caps.h`,
`cuda_arch_tactics.cu`, `cuda_paged_attn.cu` and `tests/vt/test_ops_attention.cpp`
are byte-identical to `main`, so there is nothing left to conflict.

The measurement survives the revert, and that is arithmetic rather than
assertion: the binary that produced 7.680 s carried the opt-in call, but the
helper returns immediately below 49,152 B, so it was a no-op on every launch
those numbers came from. The withdrawn evidence is the `test_ops_attention` run,
which measured code no longer in the tree; it is marked WITHDRAWN in
`.agents/benchmark-record.md` rather than deleted.

One consequence is disclosed rather than left to be found: the f32 L2 parity arm
at production geometry now refuses instead of running slowly. It fails loud in
both worlds -- a `cudaGetLastError` throw at `cuda_ops.cu:3352` today, a
`VT_CHECK` naming the head_dim once #1578 lands -- and nothing gated reaches it,
since production is bf16 and the f32 arm is exercised at the fixture's reduced
dimensions. Filed under `## Owed` against #1612.

Also in this commit, because they are the same edit to the same records:

- The corrected ratio reads **~6.0x** everywhere. 6.23x survives only as the
  uncorrected upper end of the 6.03-6.23x range, never on its own.
- Two anchors were wrong at this row's declared base `6b48edb2c` as well as at
  HEAD. `ltx2_video.cpp:4055-4059` pointed at prose about the res_2s step
  counter -- and it is the citation the whole reachability argument rests on; the
  production call site is `:4246`. `cpu_ops.cpp:3551-3562` pointed at
  `FusedStore`; the three CPU attention registrations are at `:3750-3761`. Every
  other anchor in the spec was re-read against the base rather than carried
  forward, and eight more were corrected.
- The committed A/B harness had a third precondition that grepped for
  `FlashTileSmemOptIn`, a spelling no revision of this change ever used, so it
  counted 0 and would have `exit 42`-ed on a correct tree. It is removed with the
  cap-raise it guarded, and the comment says why a precondition that cannot pass
  is not a stricter one.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
merge(row/LTX25-DIT-ATTN-FLASH): take origin/main c020347 under the b…
416bcb5 
…ranch

Clean automatic merge. `.agents/issue-index.md` union-merged at `4 0` -- four
appended rows, no deletion, no duplicated id -- and this branch's diff against
the merged main is confined to the ten files the row owns. In particular nothing
under `src/vt/cuda/` or `tests/vt/` appears in it, which is the check that the
reverted shared-memory cap-raise left no residue behind to conflict with #1578.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
docs(#1549): refresh the four anchors this change's own comment moved
382414e 
Adding the bf16-tile arithmetic comment to `AttentionDev` shifted the call sites
below it, and the stream-dtype `VT_CHECK` with them. The spec now gives both
numbers where they differ -- the base `6b48edb2c` line and the post-change line --
rather than one that is right at neither revision, which is the failure the two
corrected anchors in the previous commit already were.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
docs(#1549): bring the LTX attention row back inside the table budgets
47e9abb 
`check-public-doc-tables.py` reds a cell over 220 chars and a row over 600. The
previous commit's addition to the caveat column took that row to 408 and 680, so
the forensics move to `.agents/benchmark-record.md` -- which is where the checker
says they belong, and where the withdrawn evidence is already written out in
full. The row still names all four things a reader of the ratio needs: the
`PENDING` A/B, the four ways the arms differ, the missing pixel gate, and that
the cap-raise is reverted with its head_dim evidence withdrawn.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@mudler
docs(#1549): say exactly how the f32 arm refuses, in each of the two …
dc979aa 
…worlds

The comment said "refuses, loudly and by name". Only half of that is true today:
`Check(cudaGetLastError(), "attention-dense-flash launch")` at `cuda_ops.cu:3352`
names the OP but not the head_dim, and it is #1578's `VT_CHECK` that will name
the head_dim too. Both are loud and neither is silent, which is the property that
makes this a disclosure rather than a blocker -- so the comment now states the
two cases separately instead of claiming the stronger one for both.

Anchors resynced to the lines this edit moved.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
localai-bot added a commit that referenced this pull request Aug 22, 2026
@localai-bot @mudler
fix(KERNEL-ATTN-DENSE-FLASH): a model on the naive attention kernel n…
063133e 
…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>
@mudler
merge(LTX25-DIT-ATTN-FLASH): take origin/main 063133e, and re-append …
0910c2b 
…this row's record sections at the file's TRUE end

Two record files collided, and both are the shape AGENTS.md `## Records` names as a
lock: a shared surface that every pull request writes. Neither collision is a
disagreement about a fact, so neither is resolved by choosing a side.

`.agents/benchmark-record.md` is a tail-append collision. This row appended its
`LTX25-DIT-ATTN-FLASH` section and `origin/main` appended `SPEC-DFLASH2 W6` at the
same place. Both are kept: main's section stays where main put it and this row's 173
lines are re-appended after it, at the file's true end. Against `origin/main` the
file reads `173 0` in `git diff --numstat` with zero deleted lines, so nothing of
main's was traded away to make room.

`docs/STATUS.md` is a keyed record and was resolved as one, by taking the complete
target-branch row and applying this row's scoped edit again rather than accepting
either side whole. Two rows conflicted. `Speculative decoding` is main's alone -- this
branch never touched it, verified byte-for-byte against the merge base `c020347a7` --
so main's text is taken unchanged. `Image, video, audio, speech, music, and diffusion
models` was edited by BOTH: main rewrote the MiniMax-Music3 clause (595.9 s to
449.969 s, the DiT falling from 62.4 % to 50.2 %, citation gaining #1555) while this
row appends the LTX-2.5 DiT sentence after it. Composing main's row with this row's
addition leaves `docs/STATUS.md` reading `1 1` against `origin/main`: exactly the one
row this change owns, with every other row byte-for-byte main's.

The #1578 interaction was re-measured here rather than assumed, because this row
removes two more allowlisted `vt::Attention` sites. 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 `ltx2` and
`ltx2_device` as `STALE (not a failure)` -- the outcome
`scripts/attention-rung-allowlist.txt` documents for a removing row, since deleting
the stems here would red `test_allowlist_holds_only_the_in_flight_stems` unless its
expected set moved in the same change.

`.agents/issue-index.md` unioned cleanly: `4 0` against `origin/main`, no issue id
twice in the file.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]
@localai-bot
localai-bot merged commit 90e8c3c into main Aug 22, 2026
22 of 27 checks passed
This was referenced Aug 22, 2026
Closed
Merged
localai-bot added a commit that referenced this pull request Aug 22, 2026
@localai-bot @mudler
fix(#1663): a stale attention-rung stem excused a DELETED marker in t…
580e854 
…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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

LTX-2.5 DiT self-attention runs the correctness-grade kernel: 47.84 s per forward, and no caller is told

2 participants

@localai-bot @mudler