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

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