fix(SPEC-DFLASH2): drop the unistd.h include #1581 had already made redundant (#1617) by localai-bot · Pull Request #1618 · mudler/vllm.cpp

localai-bot · 2026-08-21T17:23:58Z

d27639e71 added #include <unistd.h> to test_qwen3_dflash2_gguf.cpp to
declare a ::getpid() call that 13548db8f had already removed, three commits
earlier, by routing the same need through vllm_test::ProcessId() and
support/process_id.h. The change was authored against the pre-#1581 tree,
where it was the correct fix, and landed after the call it supported was gone.

The file now carries no ::getpid call at all, so the include declares nothing
it uses. It is also the one spelling tests/support/process_id.h exists to
prevent: MSVC ships no <unistd.h>, so an unguarded include does not merely
warn on Windows, it fails to compile, which is the #603 shape of the #503 class.
It did not surface as a new red because both windows-msvc-* lanes are
baseline-red and never run on main.

Verified by compiling the translation unit with the project's own flags,
including -Werror, after the removal: rc 0, object produced. The file still
includes support/process_id.h and still calls vllm_test::ProcessId().

Found while merging external contributor pull requests and fixed in the same
flow; it is not that work's defect.

Closes #1617

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-opus-5 [Claude Code]

…edundant (#1617) `d27639e71` added `#include <unistd.h>` to test_qwen3_dflash2_gguf.cpp to declare a `::getpid()` call that `13548db8f` had already removed, three commits earlier, by routing the same need through `vllm_test::ProcessId()` and `support/process_id.h`. The change was authored against the pre-#1581 tree, where it was the correct fix, and landed after the call it supported was gone. The file now carries no `::getpid` call at all, so the include declares nothing it uses. It is also the one spelling `tests/support/process_id.h` exists to prevent: MSVC ships no `<unistd.h>`, so an unguarded include does not merely warn on Windows, it fails to compile, which is the #603 shape of the #503 class. It did not surface as a new red because both `windows-msvc-*` lanes are baseline-red and never run on `main`. Verified by compiling the translation unit with the project's own flags, including -Werror, after the removal: rc 0, object produced. The file still includes `support/process_id.h` and still calls `vllm_test::ProcessId()`. Found while merging external contributor pull requests and fixed in the same flow; it is not that work's defect. Closes #1617 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code]

…he redundant include this branch had also dropped `main` moved twice while this branch was being reconciled. `13548db8f` (#1581) landed the `tests/support/process_id.h` seam and `369a0c2aa` (#1618) then dropped the bare `#include <unistd.h>` that #1498 had added for the same defect, so the file is back to one repair. This branch already held `main`'s version of it whole, and `git diff origin/main -- tests/vllm/models/test_qwen3_dflash2_gguf.cpp` is still empty after this merge. The merge is taken so `check-commit-style` and `check-commit-trailers` can RUN: both refuse a range whose base is not an ancestor of its head, and a moved `origin/main` turns them into SKIPs rather than passes. `.agents/issue-index.md` was the only file both sides touched and merged as a pure union-append, checked by multiset against both parents -- 555 and 549 lines in, 557 out, no parent line dropped, altered or invented. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

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

… the correctness-grade attention kernel (#1545) (#1579) `src/vllm/model_executor/models/muse_glimmer_vision.cpp:639` called `vt::Attention` for every segment of every one of the perception encoder's 50 blocks. `src/vt/ops.cpp:2680` resolves that op to the kernel whose own header at `src/vt/cuda/cuda_ops.cu:1456-1460` calls itself "Correctness-grade (M0.9)": one 256-thread block per (query, head), a 256-wide shared-memory tree reduction for every key, and no K/V tiling. The tower has no second attention path, no knob and no A/B rung, so nothing could route around it. It now names `vt::AttentionDenseFlash`, the rung `whisper_audio.cpp:310-322` and `qwen3_vl_vision.cpp:462-480` already default to. `AttentionDenseFa2` is not usable at head_dim 96 and would fall through to the same kernel anyway. The size of the problem is now read rather than inferred. #1545 labelled its token count illustrative and asked whoever took it to pin the shipped resolution first, because it moves the estimate quadratically, and the released checkpoint was already on the NAS. `config.json` ships `layer_types` explicitly -- 13 `full_attention` and 37 `window_attention`, exactly what `muse_glimmer_weights.cpp:429-439` derives when the key is absent -- and `pos_emb_height/width` put the window at 32x32 = 1024 patches, so a window layer does not attend over the whole image. `processor_config.json` sets `max_image_tokens` to 4096. Whether that counts patch tokens or post-merge tokens is not decidable from the checkpoint, so both readings are carried: 34 s or 375 s of naive attention per image, against the 4.8 s the issue illustrated with. This lands UNREACHED, deliberately, under the staged-slice exception in `.agents/reachability.md`. What is not reached: `MuseGlimmerVisionForward` and its two callers `MuseGlimmerEncodePixelGroups` and `MuseGlimmerGenerateGreedyViaRegistry` have no caller in `src/`, in `examples/` or in `include/vllm.h`, only in `tests/`, and `muse_glimmer_registry.cpp:13-14` says why: "The perception encoder is still W3, so an image or video prompt is a pending brick." The row that owns the wiring is `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation`, at W4 and W5 of `.agents/specs/muse-glimmer.md` §3. The issue that tracks it is #1566, filed by this change because the model's umbrella #268 is closed and nothing open named the gap; both specs list it under `## Owed`. Repairing the cost before the wiring costs one line. Repairing it afterwards costs a regression hunt through a change that did not touch attention. Numerics, both halves. On CPU the swap is byte-identical BY CONSTRUCTION: `src/vt/cpu/cpu_ops.cpp:3760` registers `kAttentionDenseFlash` to the same `AttentionKernel` function pointer `kAttention` gets at `:3750`, and each op has exactly one CPU provider. Every existing golden passes at its previous value and no tolerance was widened. On CUDA it is NOT bit-identical: the flash kernel reduces per-lane head_dim strips with `__shfl_xor_sync` under a running online softmax where the naive kernel uses a 256-wide block tree, so the f32 partial-sum grouping differs, as `include/vt/ops.h:3304-3327` states. No CUDA run backs this change -- `dgx:gpu0` is held by the operator and no lease was taken -- and the A/B is recorded as owed rather than claimed. The test had to be a ROUTING test, because a numeric one cannot see this change: on CPU both ops are the same kernel, so any comparison passes before it and after it. The new case counts op-provider selections the way `tests/vllm/models/test_ltx2.cpp:665-686` does for the same question, and asserts a DERIVED 12 -- 5 window + 2 full + 5 window calls over the fixture's segmentation -- so a partial re-route fails it too. `src/vt/cuda/cuda_ops.cu` is deliberately untouched. #1544 item 2 owes a repair to `LaunchAttentionDenseFlash`'s advertised `head_dim <= 256` contract and the LTX-2.5 row is implementing it. head_dim 96 does not need it: the production bf16 arm asks 24 KiB of the 48 KiB default cap, and the spec records that the f32 arm sits exactly on the cap and fails loud rather than silently if it ever runs on CUDA. ## How to verify ```sh cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release cmake --build build --target test_muse_glimmer_vision test_muse_glimmer_wiring -j 8 ./build/tests/test_muse_glimmer_vision ./build/tests/test_muse_glimmer_wiring ``` Observed on this branch: 8/8 cases and 103 assertions, then 9/9 cases and 10,317 assertions, both `SUCCESS!`. The routing case prints `attention selections: dense-flash +12, naive +0`. Red first, with the test in place and the production line still on `vt::Attention`: `dense-flash +0, naive +12`, and both `CHECK`s fail as `CHECK( 0 == 12 )` and `CHECK( 12 == 0 )`. Byte-identity measured, not only argued. Both binaries were built from one tree with only that line differing, and every `MESSAGE:` line of the suite was captured from each and diffed. Exactly one line of fourteen differs, and it is the routing counter. The f32 stage numbers stay at `rel_l2` 1.201e-07 to 2.983e-07 and the bf16 tower at `rel_l2=5.951e-03 max_abs=3.675e-02` on both sides. `scripts/agent-preflight.sh` reports **All gates green.** on the head, run here. An earlier run of it on this branch failed `test_cpu_x86_llamacpp_floor`, and that failure was inherited rather than caused here: it failed the same way on the unmodified base `04f1cead6`, with `waiting for quiet: 15s busy=111% builders=0 load=67.47` and exit 4 where the case wants 2, which is the load-dependent behaviour #618 tracks. The host was at loadavg 67 to 121 then and lower at the final run, and the gate passed. THREE CI lanes are red on this pull request and none is caused here. Each was verified per-job against a main baseline, by comparing the failing ASSERTION and not the job name: - `windows-msvc-cpu` and `windows-msvc-vulkan` fail at `ERROR: CMakeLists.txt: MSVC /W4 /WX policy is negated on the C/C++ compile by /w`, before any compilation. They are red on main's own scheduled baseline too, and #584 / #965 track them. - `build-test-cpu` and both `sanitize-cpu` arms fail on ONE shared doctest case, `test_runner.cpp:1557`, which landed on main in `e2a9e035d` (#1273) and is owned by #1602 and #1608. The scheduled main baseline at `e2a9e035d` fails the identical `CHECK_THROWS_WITH_AS( make_runner(), "Block size must be a multiple of 16", std::invalid_argument )` with the same "No valid attention backend for device type 0" text, in all three jobs, with zero sanitizer findings in either sanitized arm. This branch does not touch `test_runner.cpp`. `build-newest-gcc` was red when this body was first written and is GREEN here: #1581 and #1618 landed the `::getpid` repair on main and this branch was re-merged onto `2e7f3bee7` to pick it up. `test_cpu_threadpool`, which an earlier run could not verify as inherited, PASSED on the re-run, leaving `test_runner` as the whole of the `build-test-cpu` red — 1 failing test out of 585. ## What is not verified No CUDA run. `dgx:gpu0` was held by the operator, no lease was taken, and the A/B that would confirm the speed is recorded under `## Owed` in the spec rather than claimed. The cost figures above are arithmetic from the 5.70 ns/iteration GB10 constant recorded in #1544, not a measurement of this tower. Closes #1545. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>

localai-bot merged commit 369a0c2 into main Aug 21, 2026
1 of 15 checks passed

localai-bot deleted the row/FIX-UNISTD-REDUNDANT-1617 branch August 21, 2026 17:24

Navigation Menu

fix(SPEC-DFLASH2): drop the unistd.h include #1581 had already made redundant (#1617) - #1618

fix(SPEC-DFLASH2): drop the unistd.h include #1581 had already made redundant (#1617)#1618
localai-bot merged 1 commit into
mainfrom
row/FIX-UNISTD-REDUNDANT-1617

localai-bot commented Aug 21, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

localai-bot commented Aug 21, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants