diff --git a/.agents/issue-index.md b/.agents/issue-index.md index e1f5ba9e5..2ee2087db 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -538,6 +538,8 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1548](https://github.com/mudler/vllm.cpp/issues/1548) | `ENG-RELEASE-CONTAINERS` | **The `cuda` lane built its ten-SM fat binary at `$(nproc)` on a hosted runner and the runner died under it, and no job declared `timeout-minutes`.** Run [32447481128](https://github.com/mudler/vllm.cpp/actions/runs/32447481128) died at object 512 of 787, about 35 minutes in, with `the runner has received a shutdown signal` and exit 143. Not a timeout: no `timeout-minutes` appeared anywhere in `.github/workflows/containers.yml`, so the six-hour default applied. `scripts/build-linux-accelerator-release.sh:24` sets ten device architectures, so each `.cu` is compiled ten times and one compiler process holds many times the resident set of a `cpu` or `vulkan` translation unit. MITIGATED, NOT DIAGNOSED: memory exhaustion is the leading hypothesis and it is NOT proven, because GitHub infrastructure reclamation produces the same message and the same exit code and the available logs cannot separate them. The change removes the one cause this repository controls. Parallelism is now LANE-AWARE, so the `cpu` and `vulkan` lanes are not slowed: they keep `$(nproc)` and only `cuda` takes 2. The value is measured, not guessed. `.github/workflows/ci.yml:801` already builds the SAME ten-SM fat gencode set on a hosted runner at `--parallel 2` and is green, and the 512-of-787 data point puts a halved build near two hours, which answers `.agents/specs/container-images.md:200-204` and its concern that two jobs would not finish inside a hosted runner's limits. `timeout-minutes: 300` goes on both building jobs for a separate reason: under the six-hour default a hang, a reclaimed runner and an exhausted one all report the same exit 143, so the next failure is undiagnosable. The budget LABELS it rather than policing it, and is loose because no arm64 container leg has ever built. `scripts/check-container-workflow.py` gates the SHAPE and not the number, so retuning the cap needs no checker edit, and prints the resolved cap and budget in its OK line. Spec [fix-ci-container-publish.md](specs/fix-ci-container-publish.md) | bug | | [#1541](https://github.com/mudler/vllm.cpp/issues/1541) | — | **A REFUSING length guard at the request boundary: the #1365 fix removed the quadratic cost but added no bound.** `SPEC-BPE-QUADRATIC-MERGE` took 64 KB in one pretoken from 23,620.695 ms to 7.797 ms and moved the exponent from ~2 to ~1 (`67823aee2`, [#1539](https://github.com/mudler/vllm.cpp/pull/1539)); it did not add a LIMIT, and a linear cost against a 100 MB body is a smaller problem than a quadratic one rather than the absence of one. The only bound in the stack today is httplib's `CPPHTTPLIB_PAYLOAD_MAX_LENGTH` of 100 MB, and there is still no authentication anywhere in `src/vllm/entrypoints/`. Two binding constraints, both from `.agents/specs/bpe-quadratic-merge.md` `## Defence in depth`: it must REFUSE with an error naming the limit and never truncate, because silently shortening a prompt returns a model output for text the caller did not send; and it belongs at the request boundary rather than in `src/vllm/v1/engine/input_processor.cpp::ValidatePromptLen`, which needs the token count the expensive step produces and so cannot run before it -- placing the guard there reproduces the exact ordering that made the original defect reachable. A byte or character bound is checkable before any tokenization happens, which is the point. NOT a defect #1365 leaves behind and not fixed in that flow: the implementing branch carried no recorded remote-write authority, so its `## Outcome` named the filing as owed AT LANDING and the operator filed it at the merge. Owed under `## Owed` in [bpe-quadratic-merge.md](specs/bpe-quadratic-merge.md) | bug | | [#1546](https://github.com/mudler/vllm.cpp/issues/1546) | `BENCH-CLOCK-GATE-ROUTE` | **`gpu_clock_state.compare_clock_records` bounds the cross-arm MEDIAN offset and nothing bounds the difference in EXCURSION BURDEN between the arms.** The two are independent on the 2026-08-19 Qwen3.8-27B bf16 c1 evidence: `median_offset_pct` is exactly **0.0** on all three pairings while the arms time-weighted mean-clock cost differs by **0.020 / 0.153 / 0.103** points rep for rep (ours 0.136 / 0.402 / 0.204 against vLLM 0.116 / 0.249 / 0.307, re-derived from the raw `*.samples.json` at `/mnt/nas_share/rc/q38bf16/out/`). The median cannot see the excursion population, which is exactly the part that does NOT cancel between the arms and therefore the part that transfers into the ratio. Proposed: one ADDITIVE term holding the two arms mean-clock cost within the same physics ceiling `MAX_CROSS_ARM_OFFSET_PCT` already rests on, which encodes what `7e07bbc91` measured -- a workload-generated excursion appears in BOTH arms, a GPU-state defect appears in ONE. Explicit non-goal: this must NOT become a route to re-scoring the nine discarded windows, which carry two independent refusals and stay `DISCARD`. Decided in [clock-gate-route.md](specs/clock-gate-route.md) | gap | +| [#1545](https://github.com/mudler/vllm.cpp/issues/1545) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | **Muse Glimmer's perception encoder named the correctness-grade attention op for all 50 of its layers.** `src/vllm/model_executor/models/muse_glimmer_vision.cpp:639` called `vt::Attention`, which `src/vt/ops.cpp:2680` resolves 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, no K/V tiling. Sole attention path in the tower, non-causal, no knob, 50 layers, H=16, head_dim 96. An instance of the class issue [#1544](https://github.com/mudler/vllm.cpp/issues/1544), and the worst-shaped one in the tree. FIXED HERE by naming `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. The geometry is now PINNED from the released checkpoint rather than inferred (`/mnt/nas_share/checkpoints/muse-glimmer-30b/`): `layer_types` ships 13 `full_attention` + 37 `window_attention`, the window is 32x32 = 1024 patches, and `max_image_tokens` is 4096, which puts the naive cost at 34 s or 375 s per image depending on whether that counts patch or post-merge tokens -- not the 4.8 s the issue illustrated with. Lands UNREACHED: the tower has no production caller, tracked by [#1566](https://github.com/mudler/vllm.cpp/issues/1566). Spec [muse-glimmer-vision-attn-flash.md](specs/muse-glimmer-vision-attn-flash.md) | perf | +| [#1566](https://github.com/mudler/vllm.cpp/issues/1566) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | **The Muse Glimmer perception encoder has no production caller, so every change inside the tower lands unreached.** `MuseGlimmerVisionForward` is called only from `MuseGlimmerEncodePixelGroups` and `MuseGlimmerGenerateGreedyViaRegistry`, and neither has a caller in `src/`, `examples/` or `include/vllm.h` -- only `tests/`. `src/vllm/model_executor/models/muse_glimmer_registry.cpp:13-14` says so: "The perception encoder is still W3, so an image or video prompt is a pending brick." NOT fixed in this flow: wiring an image prompt through `ModelRegistry::Forward` is W4/W5 in `.agents/specs/muse-glimmer.md` §3 and is a model port, not a one-line repair. Filed while landing [#1545](https://github.com/mudler/vllm.cpp/issues/1545), which needs an OPEN issue to name for its `.agents/reachability.md` staged-slice exception and found the model's umbrella issue [#268](https://github.com/mudler/vllm.cpp/issues/268) closed. Owed under `## Owed` in [muse-glimmer.md](specs/muse-glimmer.md) | gap | | [#1544](https://github.com/mudler/vllm.cpp/issues/1544) | `KERNEL-ATTN-DENSE-FLASH` | **`vt::Attention` is opt-in-by-name with no selector, no warning and no gate, and `AttentionDenseFlash` advertises a head_dim it cannot launch.** Nine live `vt::Attention(` call sites under `src/vllm/model_executor/models/`; the op resolves `OpId::kAttention` straight to the "correctness-grade (M0.9)" kernel and NOTHING ever routes it up, so a model whose author never heard of the fast rungs pays up to ~500x with correct output. A token gate cannot see it by construction — every rung is bit-identical or inside the bf16 envelope, so the goldens pass either way. FIXED HERE ADDITIVELY, and the freeze is deliberately NOT touched: `kAttention` stays byte-identical for text decode, and the six deliberate sites stay on it, because three are reference arms a gate compares against (`nemotron_h.cpp`, `nemotron_h_device.cpp`, `qwen3_5.cpp`) and two are the `VT_*_EAGER` rungs of a same-binary A/B (`whisper_audio.cpp`, `qwen3_vl_vision.cpp`) — rerouting any of them moves the comparison rather than the shipping kernel. New `scripts/check-attention-rung-consistency.py` refuses a model TU that names `vt::Attention` with no `// VT-ATTN-NAIVE:` reason beside the call, so the six deliberate sites now say why and a new author gets a red instead of a silent 500x. A selector that auto-routes was REJECTED for exactly the reference-arm reason; the marker is per-SITE and in-file, so the ordinary case writes no shared record. Item 2: `AttentionDenseFlash` claimed `head_dim <= 256` while asking for `2*64*d*sizeof(Tin)` bytes of dynamic shared memory with no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`, so the real ceiling was CUDA's default 48 KiB cap — 192 bf16, 96 f32 — and Kimi (192 f32) or Qwen3.5 (256) would have hit a bare launch error naming nothing. The bound now lives in `include/vt/ops.h` as pure host arithmetic a GPU-less box can execute, and the launcher refuses above it naming `vt::AttentionDenseFast`, which uses no shared memory and does serve those widths. Mirrors vLLM's own polarity: `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 it by launching. Spec [attention-rung-visibility.md](specs/attention-rung-visibility.md) | bug | | [#1573](https://github.com/mudler/vllm.cpp/issues/1573) | — | **`AttentionDenseFlash`'s repaired head_dim bound is proven by arithmetic and not by a launch.** The pure-host bound (`AttentionDenseFlashMaxHeadDim`, 192 bf16 / 96 f32) is unit-tested and mutated on a CPU box, but nothing there executes `LaunchAttentionDenseFlash`, so the CUDA case asserting that head_dim 256 in f32 REFUSES and names `vt::AttentionDenseFast` — and that head_dim 96, exactly on the cap, still runs — emits a loud PENDING message and returns. Its reachability mutation (drop the `VT_CHECK` on the bound, require the case to go RED) needs the same device. NOT fixed in the flow that filed it: `dgx:gpu0` was held by the developer for the whole row and AGENTS.md forbids reaching a fleet device outside a lease, so the honest report is PENDING on a named resource rather than a skip wearing a pass. Owed under `## Owed` in [attention-rung-visibility.md](specs/attention-rung-visibility.md) (risk R1) | gap | | [#1542](https://github.com/mudler/vllm.cpp/issues/1542) | `MODEL-MUSIC-minimax-music3-mini-max-music3-for-conditional-generation` | **MiniMax-Music3's flow-matching DiT is 62.24 % of a run and has never been profiled below the stage boundary.** `.agents/specs/minimax-music3.md` §20 (merge `f0396049d`, `rc` job `c206ec87` on an idle `thor:gpu0`, checkpoint staged to local disk) prices the developer's 20 s / 30 steps configuration: `denoise.dit_device` **370.556 s / 62.24 % / 120 calls**, against `vocoder.decode_window` 122.169, `ar.lm_decode_step` 56.174 and `ar.depth_forward` 21.099. The DiT **did not move** — 370.556 s against §15.7's 370.746 s over an identical 120 calls — while everything around it got 3.45x to 18.9x faster, so it went from ~3 % of a run to 62 % without changing. §15.2's "the GPU is not the problem, and no further DiT work will move this number" is recorded as INVERTED by §20.5 and must not be resurrected: it was true of a 4 s / 4-step clip whose wall was 69.5 % CIFS load and 23.4 % host depth decoder. `VLLM_CPP_MUSIC3_PROFILE=1` reports ONE bucket for the whole forward (`minimax_music3_speech.cpp:321`), so everything known about those 370.556 s is arithmetic on the outside of a black box: ~1.544 s per forward, ~3.33 TFLOP of block-stack GEMM at seq ~690, hence **~2.2 TFLOP/s** and 44x above the 9.66 GB weight-read floor, with no measurement of which of the GEMMs, `vt::AttentionCross`, the per-call cuBLASLt descriptor and heuristic rebuild, or the host packing owns it. **The dtype is NOT the gap and is settled rather than assumed**: at the pinned diffusers oracle `c6da9936`, `scripts/convert_minimax_music3_to_diffusers.py:267` defaults `--dtype float32` and `:208` applies it to the transformer while `:214` forces ONLY the RVQ depth decoder to bf16, `denoise.py:83` casts the condition to `transformer.dtype`, and the released `transformer/` shard header reports `F32` for all 231 tensors — so a bf16 or TF32 DiT would be a divergence from the oracle rather than a repair of one. Scoped as: intra-forward profiler SPANS behind a second opt-in (`VLLM_CPP_MUSIC3_DIT_SPANS=1`, so the default path stays byte-for-byte what §20 timed and prior tables stay comparable), a cuBLASLt ceiling probe at the DiT's own shapes on `thor:gpu0`, then the lever the evidence names with correctness first at the UNCHANGED `kDitRelTol` / `kDitAbsFloor` / `kDitMeanAbsTol` bounds. Spec §21 of [`minimax-music3.md`](specs/minimax-music3.md). Follows [#1512](https://github.com/mudler/vllm.cpp/issues/1512) | perf | diff --git a/.agents/specs/muse-glimmer-vision-attn-flash.md b/.agents/specs/muse-glimmer-vision-attn-flash.md new file mode 100644 index 000000000..81a1b6027 --- /dev/null +++ b/.agents/specs/muse-glimmer-vision-attn-flash.md @@ -0,0 +1,436 @@ +# MUSE-GLIMMER-VISION-ATTN-FLASH — the perception encoder's 50 layers all ran on the correctness-grade kernel, on a path nothing calls yet + +Row: `MUSE-GLIMMER-VISION-ATTN-FLASH`, against the model-matrix row +`MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` +([`muse-glimmer.md`](muse-glimmer.md)). +Issue: [#1545](https://github.com/mudler/vllm.cpp/issues/1545), an instance of +the class issue [#1544](https://github.com/mudler/vllm.cpp/issues/1544). +Sibling instance, on the only live victim: +row `LTX25-DIT-ATTN-FLASH` / [#1549](https://github.com/mudler/vllm.cpp/issues/1549), +whose spec `ltx25-dit-attn-flash.md` is not on `main` yet, which is why it is +named here rather than linked. + +Every bare `file:line` anchor below is read at this row's base, **`04f1cead6`**, +which is `origin/main` at the claim. + +## Now + +`DONE`. The one production line is changed, a routing test that fails without it +is in the suite, and every existing Muse Glimmer golden is byte-for-byte +unchanged because the CPU dispatch for both ops is the same function pointer. +What this row does **not** carry is a CUDA measurement: no GPU lease was taken, +`dgx:gpu0` is held, and §5 records the CUDA numerics change and §8 records the +A/B as owed to whoever wires the encoder. + +## 0. The one-sentence finding + +`src/vllm/model_executor/models/muse_glimmer_vision.cpp:639` called +`vt::Attention` for every one of the perception encoder's 50 layers, that op is +frozen on the kernel whose own header at `src/vt/cuda/cuda_ops.cu:1456-1460` +calls itself "Correctness-grade (M0.9)", and the tower has no other attention +path, no knob, and no A/B rung to escape onto. + +## 1. Why this row exists although nothing calls the code + +This is the part that governs how the change lands, so it comes before the +arithmetic. + +`MuseGlimmerVisionForward` is reached only from +`MuseGlimmerEncodePixelGroups` (`muse_glimmer_mm.cpp:191-203`) and +`MuseGlimmerGenerateGreedyViaRegistry` (`muse_glimmer_mm.cpp:276-345`). Neither +has a caller in `src/`, in `examples/`, or in `include/vllm.h`; the only callers +in the tree are `tests/vllm/models/test_muse_glimmer_wiring.cpp:724,804,805,864,883` +and `tests/vllm/models/test_muse_glimmer_vision.cpp:308,348`. The registry TU +states the same thing in prose at `muse_glimmer_registry.cpp:13-14`: *"The +perception encoder is still W3, so an image or video prompt is a pending +brick."* + +So this row cannot prove reachability from a production entry point and does not +pretend to. It is a **latent-cost repair on an unreached staged slice**, and the +argument for doing it now rather than later is that the cost is not latent for +long: the moment W4 wires an image prompt into `ModelRegistry::Forward`, this +becomes a multi-second — on the shipped geometry, a multi-*minute*, see §3 — +time-to-first-token stall that will be discovered by whoever did the wiring +rather than by whoever caused it. Fixing it before the wiring costs one line. +Fixing it after costs a regression hunt through a change that did not touch +attention. + +`.agents/reachability.md`'s staged-slice exception is taken deliberately and its +four obligations are all met: the commit body and the pull request body name what +is unreached, `## Owed` below names it, the owning row is +`MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation`, and the issue +that tracks the wiring is [#1566](https://github.com/mudler/vllm.cpp/issues/1566), +filed by this row because the model's umbrella issue +[#268](https://github.com/mudler/vllm.cpp/issues/268) is CLOSED and nothing open +tracked the gap. + +## 2. The call site, and why it maps onto the target op + +`muse_glimmer_vision.cpp:611-641` is the whole attention block. Its properties, +all read at the base: + +| property | value | anchor | +|---|---|---| +| ops called | exactly one, `vt::Attention` | `:639`, sole `vt::Attention` in the file | +| causal | never | `:632`, `AttentionArgs{scale, /*causal=*/false}` | +| shape | `[seg, nh, hd]`, `hk == hq`, square | `:606-608` + `RowSlice` at `:635-638` | +| head_dim | 96 | `hidden_size 1536 / num_attention_heads 16`, `muse_glimmer_vision.h:67-68,88` | +| layers | 50 | `muse_glimmer_vision.h:69` | +| knob | none | no `getenv` in the file | + +The block-diagonal mask is **not** a mask in this implementation. It is already +expressed as a loop over `RowSlice` segments (`:633-641`), each of which is a +separate dense non-causal call over one image (full-attention layers) or one +`pos_emb_height x pos_emb_width` window (window layers). Nothing about that +structure interacts with the kernel choice, so the swap is confined to which op +the loop body names. + +`vt::AttentionDenseFlash` (`include/vt/ops.h:3317-3318`, registered at +`src/vt/cuda/cuda_ops.cu:3839-3840`) accepts exactly this contract: rank-3 +`[T,H,D]`, `key.shape[0] == query.shape[0]`, GQA broadcast with `hq % hk == 0`, +one shared float dtype, `args.causal` honoured (`cuda_ops.cu:3277-3296`). Every +`VT_CHECK` in `src/vt/ops.cpp:3083-3106` is satisfied by the tensors already +being passed. + +`vt::AttentionDenseFa2` is **not** usable and was not considered further: its +fast path requires `head_dim == 64` (`src/vt/cuda/cuda_ops.cu:3395-3400`), and +at 96 it would fall straight through to `AttentionDenseFlash` anyway, so naming +it would only obscure which kernel runs. + +## 3. The cost, with the resolution PINNED rather than assumed + +[#1545](https://github.com/mudler/vllm.cpp/issues/1545) labelled its token count +as inferred and asked whoever took the row to pin the real resolution first, +because it moves the estimate quadratically. It was cheap to pin: the released +checkpoint is already on the NAS at +`/mnt/nas_share/checkpoints/muse-glimmer-30b/`, no download needed. + +**Read from `config.json`** (`vision_config`), and it agrees with our defaults +key for key: `hidden_size 1536`, `num_attention_heads 16` (head_dim 96), +`num_hidden_layers 50`, `patch_size 14`, `merge_size 2`, `pos_emb_height 32`, +`pos_emb_width 32`. + +**The 13/37 layer split is now READ, not inferred.** `layer_types` ships +explicitly, length 50, with `full_attention` at indices +`3, 7, 11, ..., 47, 49` — 13 full and 37 `window_attention`. That is exactly +what `muse_glimmer_weights.cpp:429-439` computes when the key is absent +("full every 4th layer AND on the last"), so the fallback rule is confirmed +correct against the shipped file rather than merely plausible. + +**The window is 32x32 = 1024 patches**, from `pos_emb_height/width` through +`MuseGlimmerVisionSparsePermutation` (`muse_glimmer_vision.cpp:355-382`). This +is load-bearing for the cost model and #1545 did not have it: a window layer +does **not** attend over the whole image, it attends within one 1024-patch +block. + +**The resolution.** `processor_config.json` sets +`image_processor.max_image_tokens = 4096`, and the checkpoint README's overview +table says "Max visual tokens per image | 4,096". Whether that counts +post-merge tokens (what the language model sees) or patch tokens (what the tower +runs on) is **not** determinable from the checkpoint alone — the tree has no +image processor for this model and `transformers` is not checked out here. Both +readings are carried: + +| reading | patch tokens | grid | windows | naive cost per tower forward | +|---|---:|---|---:|---:| +| A: 4096 post-merge (merge 2, so x4) | 16,384 | 128x128 | 16 | **375 s** | +| B: 4096 patch tokens | 4,096 | 64x64 | 4 | **34 s** | +| #1545's illustrative 448x448 | 1,024 | 32x32 | 1 | 4.8 s | + +Reading A is the more likely one, and the architecture is the argument: +`window_attention` over a 32x32 block does nothing at all unless the grid +exceeds 32x32. At 448x448 the grid *is* one window, so all 50 layers compute the +identical thing and 37 of the checkpoint's layer-type entries are inert. A tower +that ships 37 window layers is a tower meant to run at grids well past one +window. + +Arithmetic, at `Q * H * K_eff` iterations and the 5.70 ns/iteration GB10 +constant stated in [#1544](https://github.com/mudler/vllm.cpp/issues/1544) and +derived there from `.agents/specs/multimodal-speed.md:24-26`, which records 56 ms +per block over 784 patches and 16 heads rather than the constant itself +(`56e-3 / (784 * 16 * 784) = 5.694e-9`). The head_dim 64 and 72 provenance is in +that issue too. The tree holds the second anchor's underlying measurement at +`.agents/benchmark-record.md:4654` — the Voxtral Whisper encoder forward over +1500 frames and 32 layers at 8870 ms — and only the 8.21 s prediction and the +"within 7%" comparison are the issue's alone: + + reading A: 13 x (16384 x 16 x 16384) = 5.583e10 + + 37 x (16 x (1024 x 16 x 1024)) = 9.932e9 + 6.576e10 x 5.70 ns = 375 s + + reading B: 13 x (4096 x 16 x 4096) = 3.489e9 + + 37 x (4 x (1024 x 16 x 1024)) = 2.483e9 + 5.973e9 x 5.70 ns = 34.0 s + +The 5.70 ns constant is anchored at head_dim 64 and 72 and its transfer to 96 is +**inferred**, exactly as #1544 labels it. Nothing in this row depends on the +constant being right to better than a factor: at reading B the naive kernel +already costs half a minute per image, and at reading A it costs six minutes. + +## 4. What changes + +One line, `muse_glimmer_vision.cpp:639`: + +``` +- vt::Attention(q, os, qs, ks, vs, aargs); ++ vt::AttentionDenseFlash(q, os, qs, ks, vs, aargs); +``` + +plus the comment above it that says why, mirroring the shape of +`qwen3_vl_vision.cpp:462-480` and `whisper_audio.cpp:310-322`, which +[#1544](https://github.com/mudler/vllm.cpp/issues/1544) names as the intended +pattern. + +**No A/B knob is added, and that is deliberate.** Qwen3-VL and Whisper each keep +a `VT_*_EAGER` escape to the naive arm because each has a measured result to +defend and a live path to measure on. This tower has neither: it is unreached, so +a knob here would add a second dead arm and a second untested branch to a file +that already has no caller. When W4 wires the encoder and someone takes the A/B +in §8, adding the knob is a two-line change made *with* the measurement in hand. + +**`src/vt/cuda/cuda_ops.cu` is deliberately NOT touched.** +[#1544](https://github.com/mudler/vllm.cpp/issues/1544) item 2 owes a repair to +`LaunchAttentionDenseFlash`'s advertised `head_dim <= 256` contract, and +row `LTX25-DIT-ATTN-FLASH` ([#1549](https://github.com/mudler/vllm.cpp/issues/1549)) +§4.3 is implementing it. +Two rows editing one kernel launcher is the shared-file lock this repository's +records rules exist to avoid. §6 records why 96 does not need that repair to be +correct today. + +## 5. Numerics: what is bit-identical and what is not + +This has two answers and they must not be collapsed into one. + +**On CPU the swap is byte-identical, by construction and not by tolerance.** +`src/vt/cpu/cpu_ops.cpp:3760-3761` registers `OpId::kAttentionDenseFlash` on +`DeviceType::kCPU` to `&AttentionKernel` — the *same function pointer* +`OpId::kAttention` is registered to at `:3750-3751`. The comment there says so: +"Flash-tiled dense attention is a CUDA shared-memory optimization; +byte-identical to kAttention on CPU." Every Muse Glimmer golden in the tree is a +CPU golden, so all of them are unchanged bit-for-bit, and that is what the gate +run in §7 observes. + +**On CUDA the swap is NOT bit-identical, and no CUDA run was taken here.** +`AttentionDenseFlashKernel` (`cuda_ops.cu:3239-3325`) is one warp per +(query, head) with the head_dim split into per-lane strips reduced by +`__shfl_xor_sync`, carrying a running online-softmax `(m, l, acc)`. +`AttentionKernel` (`cuda_ops.cu:1463`) is a 256-thread block per (query, head) +with a shared-memory tree reduction per key. The arithmetic is the same f32 +softmax; the **partial-sum grouping and the reduction order are not**, so the +two differ within the f32 envelope. `include/vt/ops.h:3304-3318` states exactly +this and says adoption is per token-exact gate. + +Three things follow, and this row states all of them rather than the convenient +one: + +1. The change **does** move CUDA numerics on this path. +2. It moves them onto the rung Whisper's encoder and the Qwen3-VL tower already + run on by default, which is where the tree's other non-causal encoder towers + already are. +3. It is bit-identical to `AttentionDenseFast` (`ops.h:3313-3314`: "the CUDA + output is BIT-IDENTICAL to it (K/V bytes merely sourced from shared memory)"). + +No tolerance anywhere was widened to accommodate any of this. The existing +`rel_l2 < 1e-6` f32 bound and `rel_l2 < 2e-2` bf16 bound in +`tests/vllm/models/test_muse_glimmer_vision.cpp` are untouched and still pass at +the same numbers, because on CPU nothing moved at all. + +The CUDA-side confirmation is **owed**, not claimed. It is recorded in `## Owed` +and in §8, and it is not fabricated as a pass. + +## 6. Shared memory at head_dim 96, for both arms + +`LaunchAttentionDenseFlash` requests `2 * kFlashBc * d * sizeof(Tin)` bytes of +dynamic shared memory (`cuda_ops.cu:3338`) with `kFlashBc = 64` +(`cuda_ops.cu:3236`), and `cuda_ops.cu` contains no `cudaFuncSetAttribute`, so +the real cap is the 48 KiB every architecture gives without opting in — the +defect [#1544](https://github.com/mudler/vllm.cpp/issues/1544) item 2 records. + +The scope of that sentence is `cuda_ops.cu` and not `src/vt/cuda/`, deliberately. +Seven files under `src/vt/cuda/` do call `cudaFuncSetAttribute` — `cuda_gdn.cu`, +`cuda_marlin_repack.cu`, `cuda_mla_attn.cu`, `cuda_paged_attn.cu`, the two other +Marlin translation units, and `flash_attn/src/flash_fwd_launch_template.h`, which +is inside the very subtree `ops.h:3309-3310` says this kernel was ported from. +None of them is on this launch path, and the claim the argument needs is the +narrow one. + +The kernel body also declares **no static `__shared__`** (`cuda_ops.cu:3239-3328` +holds one `extern __shared__` array and nothing else), so nothing competes with +the dynamic request. That check matters: one static byte would put the f32 arm +over the cap. + +| arm | `sizeof(Tin)` | request | vs the 49,152 B cap | +|---|---:|---:|---| +| bf16 — the production dtype, `muse_glimmer_vision.h:86` | 2 | 24,576 B | 50% of it | +| f32 — the per-stage gate arm, `muse_glimmer_vision.h:84-86` | 4 | 49,152 B | **exactly on it** | + +The default arm has a factor of two in hand. The f32 arm lands on the boundary +to the byte. It is a `<=` comparison and 49,152 B is the guaranteed default, so +it should launch; the kernel fails loud if it does not +(`Check(cudaGetLastError(), ...)` at `cuda_ops.cu:3352`), so the failure mode is +a diagnosable refusal and never a silent wrong answer. That claim is **not +measured here** — no GPU lease — and the f32 arm has never run on CUDA in this +tree in any case, since every f32 Muse Glimmer test is a CPU test. The +`LTX25-DIT-ATTN-FLASH` ([#1549](https://github.com/mudler/vllm.cpp/issues/1549)) +§4.3 repair removes the +question entirely by opting in to the larger cap, and this row is a beneficiary +of it rather than a blocker on it. + +## 7. Tests and gates + +**The problem this row had to solve first:** no CPU test can detect the routing +change through its *output*, because §5's first paragraph makes the two ops the +same function on CPU. A test that compared numbers would pass before the change +and after it, which is a test that measures nothing +([[a-mutation-that-never-applied-reads-as-a-passing-test]]). + +The instrument that does see it is the op-provider selection counter, and the +tree already uses it for exactly this question at +`tests/vllm/models/test_ltx2.cpp:665-686`. New case +`muse_glimmer_vision_tower_routes_attention_to_dense_flash` in +`tests/vllm/models/test_muse_glimmer_vision.cpp`: + +- `vt::EnableOpProviderCallStats(true)`, read `selections` for + `kAttentionDenseFlash` and `kAttention` on `kCPU` before and after one + `MuseGlimmerVisionForward` over the existing fixture; +- assert `kAttentionDenseFlash` gained **exactly 12** and `kAttention` gained + **exactly 0**. + +12 is derived, not observed: the fixture is 3 layers `{window, full, window}` +over two images whose grids are 6x6 and 4x4 with a 4x4 window +(`test_muse_glimmer_vision.cpp:64-78,122-134`), so the window segmentation is +`[16, 8, 8, 4]` + `[16]` = 5 calls and the full segmentation is `[36]` + `[16]` += 2 calls: `5 + 2 + 5 = 12`. Asserting the derived count rather than a +before/after inequality is what makes the case sensitive to a partial routing +change, and it pins the segmentation as a side effect. + +**Red-first evidence, observed.** With the case in the suite and +`muse_glimmer_vision.cpp:639` still on `vt::Attention`: + + :395: MESSAGE: attention selections: dense-flash +0, naive +12 + :397: ERROR: CHECK( flash_after - flash_before == 12ull ) values: CHECK( 0 == 12 ) + :398: ERROR: CHECK( naive_after - naive_before == 0ull ) values: CHECK( 12 == 0 ) + [doctest] assertions: 5 | 3 passed | 2 failed | Status: FAILURE! + +The exact inversion the derivation predicts, and 12 is confirmed by the run +rather than only by the arithmetic. After the one-line change, the whole file: + + :395: MESSAGE: attention selections: dense-flash +12, naive +0 + [doctest] test cases: 8 | 8 passed | 0 failed | 0 skipped + [doctest] assertions: 103 | 103 passed | 0 failed | Status: SUCCESS! + +`tests/vllm/models/test_muse_glimmer_wiring.cpp` is green on the same tree at +9/9 cases and 10,317 assertions. + +**The byte-identity claim of §5 is 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: + + 13c13 + < attention selections: dense-flash +0, naive +12 + > attention selections: dense-flash +12, naive +0 + +The other thirteen — `ln_pre` and `block 0` and the tower and the adapter in f32 +at `rel_l2` 1.201e-07 to 2.983e-07, and the tower in bf16 at +`rel_l2=5.951e-03 max_abs=3.675e-02` — are byte-for-byte equal across the two +binaries. That is agreement to the four significant digits doctest prints, on top +of the structural argument in §5 that the CPU dispatch is one function pointer. + +**Gate:** `scripts/agent-preflight.sh`, run by this row on its own tree. CPU +only. No GPU axis is claimed, opened, or needed for what this row asserts. + +**No public document changes, and that is the correct answer rather than an +omission.** AGENTS.md's projection table triggers on a lifecycle change, a +measurement, or a user-visible surface. This row changes none of them: it takes +no measurement (§8), it moves no roadmap row's state, and the surface it touches +is not user-visible because nothing reaches it (§1). `docs/FEATURES.md:150` +already records the vision arm's real state — "vision: **no reference run of any +kind**" — and a line saying the unreachable tower now names a faster op would be +a claim about a capability no user can invoke. + +## 8. The A/B, which is OWED and not taken + +The confirming measurement — naive vs flash on the real geometry, same binary, +`dgx:gpu0` — is not takeable by this row: the box is held by the operator and +this row has no lease authority. It is also not *worth* taking until W4 lands, +because the tower has no production caller to measure through and a synthetic +harness would measure a kernel rather than a capability. + +When it is taken it needs: the real image resolution resolved between §3's +readings A and B, both arms from one binary, and the tower forward timed end to +end rather than the kernel alone +([[component-speedup-is-not-system-speedup-fixed-serial-term]]). The recorded +precedent for the shape of the win is the Qwen3-VL tower at 14.3x +(2114 -> 148 ms/image, `bbaa182b6`), on a different geometry; it is cited as +precedent and is **not** this row's number. + +## 9. Risks + +| risk | handling | +|---|---| +| CUDA numerics move on an ungated path | Stated in §5 rather than hidden. Bit-identical to `AttentionDenseFast`, and the rung Whisper and Qwen3-VL default to. Confirmation owed. | +| f32 CUDA arm sits exactly on the 48 KiB cap | §6. Fails loud, never silent. Removed outright by the #1544 item-2 repair the sibling row carries. | +| The routing test's `12` goes stale if the fixture changes | It is derived in a comment beside the assertion, so a fixture change that breaks it says which number to re-derive and why. | +| Landing unreached code | §1. The staged-slice exception is taken explicitly with all four obligations met. | + +## 10. Stop conditions + +Stop and return `NEEDS_DECISION` if `AttentionDenseFlash` turns out not to +express the segmentation (it does — §2), if the swap moves a CPU golden (it +cannot — §5), or if a production caller for the tower appears (none exists — +§1). None of the three fired. + +## Owed + +- [#1566](https://github.com/mudler/vllm.cpp/issues/1566) — **the perception + encoder has no production caller.** `MuseGlimmerEncodePixelGroups` and + `MuseGlimmerGenerateGreedyViaRegistry` are reachable only from `tests/`, so an + image or video prompt through `ModelRegistry::Forward` still hits the W3/W4 + brick at `muse_glimmer_registry.cpp:13-14`. This row's change lands *inside* + that unreached slice and is the staged-slice exception in + `.agents/reachability.md`. Owned by + `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation`. Filed by this + row because [#268](https://github.com/mudler/vllm.cpp/issues/268) is closed and + nothing open tracked the gap. +- **The CUDA A/B is not taken** (§8). `dgx:gpu0` was held by the operator for + the whole of this row and no lease was requested. Ready-to-measure, not + measured: the recipe is in §8 and the geometry question it needs answered is + in §3. +- **The f32 CUDA arm has never launched at 49,152 B** (§6). The bf16 default arm + has a factor of two in hand; the f32 arm sits on the 48 KiB cap to the byte, and + §6 argues it fits rather than measuring it. The argument is an inference about + driver behaviour at the exact boundary and it does not account for any per-block + driver shared-memory reservation. It is safe to carry today only because the arm + is dead twice over — no production caller for the tower at all, and every Muse + Glimmer f32 test is a CPU test — and because the failure mode is a loud + `Check(cudaGetLastError(), ...)`. It stops being safe the moment W4 wires the + encoder or anyone runs the f32 stage gate on CUDA, whichever comes first. Owed: + one launch probe at head_dim 96 f32, or the `LTX25-DIT-ATTN-FLASH` + ([#1549](https://github.com/mudler/vllm.cpp/issues/1549)) §4.3 opt-in landing + first and removing the question. Recorded as debt rather than as a settled + argument, because an inference quoted twice starts reading as a measurement. +- **The image resolution is still two readings, not one** (§3). Resolving it + needs the upstream `MuseGlimmerImageProcessor`, which is neither in this tree + nor in a local `transformers` checkout. It changes the size of the number and + nothing about the correctness of the change. + +## Outcome + +Measured: nothing on CUDA, deliberately, and §8 says why. What was *established* +is the shape of the cost with the geometry pinned from the released checkpoint +rather than assumed (§3), which moved the headline from #1545's illustrative +4.8 s to 34 s or 375 s per image depending on a reading that is still open. + +Rejected: `AttentionDenseFa2`, which refuses head_dim 96 and would silently be +`AttentionDenseFlash` anyway (§2). An A/B knob, because a second arm on a path +with no first caller is two dead arms (§4). Touching +`LaunchAttentionDenseFlash`'s shared-memory opt-in, because a sibling row is +already there and one kernel launcher edited by two rows is the exact +shared-file lock the records rules forbid (§4). + +Why the default has its value: `AttentionDenseFlash` and not `AttentionDenseFast` +because the tower's whole problem is redundant global K/V re-reads over a +1024-to-16384-token non-causal context, which is the difference between those two +rungs; and unconditional rather than shape-gated because every shape this call +site can produce is inside the op's contract (§2). diff --git a/.agents/specs/muse-glimmer.md b/.agents/specs/muse-glimmer.md index 4a94aa27e..74be59a5b 100644 --- a/.agents/specs/muse-glimmer.md +++ b/.agents/specs/muse-glimmer.md @@ -1190,3 +1190,21 @@ about magnitudes. **Two greedy prompts are not a token-exactness claim.** other member of the same class in the same case (`bdiff <= 1e-5`) and deliberately leaves this one to its own derivation. A repair owes a measured floor or a precision argument, never a bigger constant. + +- [#1566](https://github.com/mudler/vllm.cpp/issues/1566) — **the perception + encoder is an UNREACHED staged slice, so everything W3 built is dead code + today.** `MuseGlimmerVisionForward` is reached only from + `MuseGlimmerEncodePixelGroups` (`muse_glimmer_mm.cpp:191`) and + `MuseGlimmerGenerateGreedyViaRegistry` (`:276`), and neither has a caller in + `src/`, in `examples/` or in `include/vllm.h` — only in `tests/`. + `muse_glimmer_registry.cpp:13-14` states it: *"The perception encoder is still + W3, so an image or video prompt is a pending brick."* §3's W4 and W5 own the + wiring. Until they land, every change inside the tower is the staged-slice + exception in `.agents/reachability.md` and has to name this issue in its commit + body and its pull request body. The first one that does is + [#1545](https://github.com/mudler/vllm.cpp/issues/1545), which routes the + tower's 50 attention layers off the correctness-grade kernel + ([muse-glimmer-vision-attn-flash.md](muse-glimmer-vision-attn-flash.md)). The + issue exists because the model's umbrella + [#268](https://github.com/mudler/vllm.cpp/issues/268) is closed, so there was + nothing open left to name. diff --git a/src/vllm/model_executor/models/muse_glimmer_vision.cpp b/src/vllm/model_executor/models/muse_glimmer_vision.cpp index 73c842732..4e4bb4da4 100644 --- a/src/vllm/model_executor/models/muse_glimmer_vision.cpp +++ b/src/vllm/model_executor/models/muse_glimmer_vision.cpp @@ -16,7 +16,7 @@ // is_neox_style=True, fp32 compute). // // Composed from the public vt:: ops (MatmulBT / Add / LayerNorm / GeluErf / -// RopeFromCache / Attention / IndexSelect) and the shared merged-QKV seam +// RopeFromCache / AttentionDenseFlash / IndexSelect) and the shared merged-QKV seam // vllm::models::FusedMergedQkvBiasSplit — no hand-rolled parallel path. The // host precomputes (patchify, positional interpolation, the 2D-RoPE cos|sin // table, the window permutation, the pixel shuffle) are deterministic f32, as @@ -630,13 +630,29 @@ std::vector MuseGlimmerVisionForward(const std::vector(n); Tensor qs = RowSlice(q3, off, seg); const Tensor ks = RowSlice(k3, off, seg); const Tensor vs = RowSlice(v3, off, seg); Tensor os = RowSlice(ao.tensor(), off, seg); - vt::Attention(q, os, qs, ks, vs, aargs); + vt::AttentionDenseFlash(q, os, qs, ks, vs, aargs); off += seg; } } diff --git a/tests/vllm/models/test_muse_glimmer_vision.cpp b/tests/vllm/models/test_muse_glimmer_vision.cpp index 4f7b80389..77ad1da57 100644 --- a/tests/vllm/models/test_muse_glimmer_vision.cpp +++ b/tests/vllm/models/test_muse_glimmer_vision.cpp @@ -29,6 +29,8 @@ #include "vllm/model_executor/models/muse_glimmer_vision.h" #include "vt/backend.h" #include "vt/dtype.h" +#include "vt/op_provider.h" +#include "vt/ops.h" #include "muse_glimmer_vision_goldens.inc" @@ -351,3 +353,55 @@ TEST_CASE("muse_glimmer_vision_tower_bf16_within_envelope") { MESSAGE("tower output (bf16): ", Fmt(et)); CHECK_MESSAGE(et.rel_l2 < 2e-2, "bf16 tower output: ", Fmt(et)); } + +// --- ROUTING: WHICH attention op the tower names ----------------------------- +// The 50 blocks are the whole cost of this tower, and it must not pay them on +// the "Correctness-grade (M0.9)" kernel `vt::Attention` is frozen on +// (src/vt/cuda/cuda_ops.cu:1456-1460) -- issue #1545, class issue #1544. +// +// No case above can see that choice, BY CONSTRUCTION rather than by oversight: +// on CPU `kAttentionDenseFlash` is registered to the SAME `AttentionKernel` +// function pointer as `kAttention` (src/vt/cpu/cpu_ops.cpp:3750-3761), so every +// number in this file is byte-identical whichever op the forward names. A +// numeric assertion would therefore pass before the routing change and after +// it, which is an assertion that measures nothing. The op-provider SELECTION +// counter is the instrument that does see it, used here exactly as +// tests/vllm/models/test_ltx2.cpp:665-686 uses it for the same question. +// +// The expected count is DERIVED, so this fails on a PARTIAL re-route and not +// only on a total one. FixtureConfig is 3 layers {window, full, window} over +// FixtureImages' 12x12 (grid 6x6) and 8x8 (grid 4x4) images with a 4x4 window +// block, so a window layer segments as [16,8,8,4] + [16] = 5 calls and a full +// layer as [36] + [16] = 2: 5 + 2 + 5 = 12. +TEST_CASE("muse_glimmer_vision_tower_routes_attention_to_dense_flash") { + vt::Backend* cpu = vt::TryGetBackend(vt::DeviceType::kCPU); + REQUIRE(cpu != nullptr); + const MuseGlimmerVisionConfig cfg = FixtureConfig(); + const MuseGlimmerVisionWeights w = FixtureWeights(cfg); + + vt::EnableOpProviderCallStats(true); + const unsigned long long flash_before = + vt::GetOpProviderStats(vt::OpId::kAttentionDenseFlash, vt::DeviceType::kCPU).selections; + const unsigned long long naive_before = + vt::GetOpProviderStats(vt::OpId::kAttention, vt::DeviceType::kCPU).selections; + const std::vector tower = + vllm::multimodal::MuseGlimmerVisionForward(FixtureImages(), w, cfg, *cpu, nullptr); + const unsigned long long flash_after = + vt::GetOpProviderStats(vt::OpId::kAttentionDenseFlash, vt::DeviceType::kCPU).selections; + const unsigned long long naive_after = + vt::GetOpProviderStats(vt::OpId::kAttention, vt::DeviceType::kCPU).selections; + vt::EnableOpProviderCallStats(false); + + MESSAGE("attention selections: dense-flash +", flash_after - flash_before, ", naive +", + naive_after - naive_before); + CHECK(flash_after - flash_before == 12ull); + CHECK(naive_after - naive_before == 0ull); + + // ...and the routed forward is still the GATED forward. On CPU this is + // byte-identical to the case above for the reason in the comment; the bound is + // the same 2e-2 bf16 envelope and is NOT widened here. + const Err et = Compare(tower, muse_glimmer_vision_ref::kTowerOut, + std::size(muse_glimmer_vision_ref::kTowerOut)); + MESSAGE("tower output (routed, bf16): ", Fmt(et)); + CHECK_MESSAGE(et.rel_l2 < 2e-2, "routed tower output: ", Fmt(et)); +}