diff --git a/.agents/claims/CLAIM-ATTN-RUNG-VISIBLE.md b/.agents/claims/CLAIM-ATTN-RUNG-VISIBLE.md new file mode 100644 index 000000000..7b5d3ecf2 --- /dev/null +++ b/.agents/claims/CLAIM-ATTN-RUNG-VISIBLE.md @@ -0,0 +1,5 @@ +# CLAIM-ATTN-RUNG-VISIBLE + +| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-ATTN-RUNG-VISIBLE` | `KERNEL-ATTN-DENSE-FLASH` (`ACTIVE`) | Claude Code (opus-5), fresh implementer — review goes to a different agent | isolated worktree `.claude/worktrees/agent-ac451652cd9b2f780`; no GPU (`dgx:gpu0` held by the developer), CPU build only | `row/KERNEL-ATTN-DENSE-FLASH`, issues [#1544](https://github.com/mudler/vllm.cpp/issues/1544) and [#1573](https://github.com/mudler/vllm.cpp/issues/1573) | Owns ONLY: `.agents/specs/attention-rung-visibility.md`; `scripts/check-attention-rung-consistency.py` and `scripts/attention-rung-allowlist.txt`; `tests/scripts/test_check_attention_rung_consistency.py`; the `// VT-ATTN-NAIVE:` marker comments at the six deliberate `vt::Attention` sites; the head_dim bound in `include/vt/ops.h` and its guard in `LaunchAttentionDenseFlash`; the new head_dim contract cases in `tests/vt/test_ops_attention.cpp`; the `KERNEL-ATTN-DENSE-FLASH` evidence cell; and the preflight / CI wiring. EXCLUDES `vt::Attention`'s behaviour and `OpId::kAttention`'s resolution, which stay frozen; EXCLUDES rerouting any of the six deliberate sites; and EXCLUDES `ltx2.cpp`, `ltx2_device.cpp` and `muse_glimmer_vision.cpp`, whose naive calls are being removed by rows in flight and which are carried on the allowlist instead | `ACTIVE` | 2026-08-21 — PR [#1578](https://github.com/mudler/vllm.cpp/pull/1578) open, body verified by `scripts/agent-pr-body.py`. Checker RED-first on the unmodified tree (the six deliberate sites, at the exact lines #1544 names), GREEN after the markers; the mutation suite green; head_dim bound and its inclusive edge pinned on CPU. Contract test MUTATED red (the old `d <= 256` bound fails 6 assertions) and restored by sha256. One leg PENDING a lease (#1573): nothing on a CPU box executes the launcher's guard. Fresh scoped review returned and its findings are REPAIRED: `check-pr-size` classification (the creation-mutation stub, 31/31 red under it), the tautological `static_assert` (register blocking hoisted to `kFlashMaxPerLane`; mutating it to 4 now reads `128 == 256` where it used to read `256 == 256`), two comments that overstated what the code guarantees, the checker's undetected spellings, the unmarked-but-excused count, the allowlist/test coupling, and the stored case count in the kernel-matrix cell | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 13cd6cdb8..8f619754b 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 | +| [#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 | | [#1555](https://github.com/mudler/vllm.cpp/issues/1555) | `MODEL-MUSIC-minimax-music3-mini-max-music3-for-conditional-generation` | **`vt::AttentionCross`'s CUDA kernel does 4.0 % of the MiniMax-Music3 DiT's arithmetic in 43.9 % of its time.** Measured on an idle `thor:gpu0` (sm_110, driver 595.78, `uptime` 3.46 with 0 logins at job start), `rc` job `0f95377f-70dd-4bf8-93b5-8e44fd762713`, checkpoint staged to local disk with `SRC_BYTES == DST_BYTES == 28517617303` and `findmnt -T /tmp/ckpt` reading `overlay`, at tree `0e18f8afd` with the intra-DiT spans of [#1542](https://github.com/mudler/vllm.cpp/issues/1542). At the shipped geometry (`dit.seq_sum / 16 = 690`, window 689 latent frames, 36 blocks, inner 2048, ff 8192) `denoise.dit_device` is 25.104 s over 16 forwards and the sixteen spans sum to 25.100 s, a 99.98 % partition: **`dit.attn` 11.010 s = 43.9 % at 0.140 TFLOP per forward = 0.204 TFLOP/s**, against the four `vt::MatmulBT` GEMMs' 13.390 s = 53.3 % at 3.334 TFLOP = **3.98 TFLOP/s**, with norms, rope, SiLU, packing and readback 0.700 s = 2.8 %. **The attention kernel is 19.5x slower per flop than the GEMMs beside it, on the same tensors, in the same forward, on the same device.** At the developer's 20 s / 30 steps that is 0.6881 s x 240 forwards = **165.1 s of the 370.510 s `denoise.dit_device` bucket and 27.6 % of the whole 598.207 s run**. MECHANISM AS A HYPOTHESIS, not a measurement: `AttentionCrossFlashKernel` (`src/vt/cuda/cuda_attention_cross.cu`) gives one warp one query row and pays a **five-step `__shfl_xor_sync` butterfly per key** for 64 useful flops (`npl = 2` at `head_dim` 64), launching `grid(44,32)` of 512 threads with `bc = 64` and 32 KiB dynamic shared memory, so occupancy is bounded near 32 warps per streaming multiprocessor — about 8 per scheduler, under what hides a 40-60 cycle dependent chain; an instruction count over 22 080 warps x 690 keys predicts ~1.8 ms per layer against a measured 19.1 ms, and every CTA re-reads the whole of K and V for 17.9 GB per forward (~65 ms, ~10 %). **No `ncu` counter was read on either side and no occupancy was measured**, so the split is measured and its attribution is not. DIRECTION, NOT A DESIGN: amortize the reduction — a lane per key so the dot product is sequential in registers, or a query TILE per warp as FlashAttention-2 does. Both reorder the head-dim summation so neither is bit-identical, which the op's contract already allows (`cuda_attention_cross.cu` NUMERICS note), and admitting one requires BOTH consumers to hold their EXISTING tolerances — this DiT at `kDitRelTol` 1e-4 / `kDitAbsFloor` 5e-5 / `kDitMeanAbsTol` 5e-6 (§14.4) and LTX-2.5, which reaches the op six times per layer. **NOT a precision change**: the DiT's fp32 is upstream's resolved choice at diffusers `c6da9936` (`convert_minimax_music3_to_diffusers.py:267` defaults `--dtype float32`, `:208` applies it, `:214` forces only the RVQ depth decoder to bf16, `denoise.py:83` casts into it, and the released `transformer/` shard header is `F32` for all 231 tensors), and cuBLASLt's fp32 emulation is separately MEASURED UNAVAILABLE here (`CUBLASLT_MATMUL_DESC_EMULATION_STRATEGY` undefined in nvcc 13.0.88, `PROBE_BUILD_RC=2`). Found by [#1542](https://github.com/mudler/vllm.cpp/issues/1542); listed under `## Owed` in §21.7 of [`minimax-music3.md`](specs/minimax-music3.md) | perf | | [#1563](https://github.com/mudler/vllm.cpp/issues/1563) | `GATE-SQUASH-SEPARATOR` | **A markdown `---` horizontal rule anywhere in a pull request body silently voids the trailer block, and `check-commit-trailers.py` blames the trailers instead of the framing.** Found 2026-08-21 writing the body for PR [#1550](https://github.com/mudler/vllm.cpp/pull/1550) ([#1542](https://github.com/mudler/vllm.cpp/issues/1542)). `parsed_trailers()` shells out to git's trailer parser, and **git treats a line of exactly `---` as the start of the patch section**, so everything after the first one is not part of the message and a trailer block below it is invisible. Reproduced with no repository state: a body of `subject / prose / --- / more prose / FOLLOWING_AGENTS_PROTOCOL / the three trailers` reports `[trailers] Following-Agents-Protocol must appear exactly once` and `[attribution] AI-Assisted must appear exactly once`; `sed -i '/^---$/d'` on that same file reports `OK: commit trailer contract`, and the `---` is the only difference. **The MESSAGE is the defect, not only the behaviour**: `Following-Agents-Protocol` appears EXACTLY ONCE in the body while the checker says it must appear exactly once, so a reader counts occurrences, finds one, counts again and dumps bytes before thinking to test the parser's own framing. `_strict_errors` already computes `_paragraphs(body)[-1]` correctly as the three trailers verbatim, so the checker holds the information needed to say "the trailer paragraph is present but git could not parse it; a `---` line at line N ends the message". Worse, the neighbouring `FOLLOWING_AGENTS_PROTOCOL must appear exactly once as a separate paragraph before the trailer paragraph` check stays SILENT, so the two errors that fire both point away from the cause. **Beyond one confusing message**: the repository sets `squash_merge_commit_message = PR_BODY`, so the body IS the landed commit message, and a body carrying a `---` lands a commit whose trailers `git interpret-trailers` cannot see, on a branch that is never force-pushed. Same permanent-damage shape AGENTS.md records for the `---------` separator GitHub wrote under `COMMIT_MESSAGES`, arriving from the AUTHOR side rather than the forge side. `scripts/agent-pr-body.py --pr ` DOES catch it and caught it here before the merge; the exposure is a body never passed through that command, which AGENTS.md notes is not a gate and cannot be one because it reaches the network, while the CI guard reads the frozen `pull_request` payload and so does not re-read a body edited after the final push. NOT FIXED HERE: it changes a checker's semantics and its message, so under `## Changing the rules or a checker` it needs its own row, a red-before test and green-after evidence. Two candidate repairs, neither chosen: name the `---` line, or strip patch-section framing before parsing so a markdown rule is inert -- the second changes what the contract accepts and is the larger decision. Suggested minimum: `tests/scripts/test_check_commit_trailers.py` gains a case pinning the reproduction above | bug | @@ -557,6 +559,8 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1613](https://github.com/mudler/vllm.cpp/issues/1613) | `GATE-QWEN38-27B-FP8-BLOCK` | **The `Qwen/Qwen3.8-27B-FP8` block-wise token gate cannot be taken, because the 28.75 GiB checkpoint is not on the share.** `/mnt/nas_share/rc/ckpt/` holds `qwen3.8-27b-hf`, which is the **bf16** artifact -- no `quantization_config` key, `text_config.dtype = bfloat16` -- and `qwen3.8-q1_0`. Neither is this subject. The share has 3.4 TiB free, so the cost is AUTHORITY: `.agents/developer-preferences.md` authorizes large downloads for the `SPEC-DFLASH2` assets only. Nothing else blocks the gate, and that was not known before: a range-request audit of all 66 shard headers at revision `017b9c7a` shows every one of the 407 `F8_E4M3` tensors has `N % 128 == 0` and `K % 128 == 0`, so the sm120 complete-scale-block refusal (#1453) that makes DSV3's `kv_a_proj_with_mqa` unservable blocks NOTHING here; the ragged GDN `in_proj_a`/`in_proj_b` `[48, 5120]` are `BF16` and named in `modules_to_not_convert`; `weight_scale_inv` ships `BF16` (byte-checked via `data_offsets`, not the label) which `LoadFp8BlockRaw` already widens by value; and the per-layer `layers-.safetensors` naming already resolves through `SelectWeightFiles`. Spec `.agents/specs/gate-qwen38-27b-fp8-block.md`, parent #1189 | gap | | [#1614](https://github.com/mudler/vllm.cpp/issues/1614) | `GATE-QWEN38-27B-FP8-BLOCK` | **Three sites said `Qwen/Qwen3.8-27B-FP8` ships "~400" `modules_to_not_convert` entries, and at revision `017b9c7a` it ships 882** (882 unique, 636 outside the vision tower). The number is the evidence for an ARGUMENT -- it is why `IsFp8BlockProjection` reads the config AND the tensors instead of probing dtypes -- so being wrong by more than 2.2x invites the next reader to re-derive it. No reading of the list produces ~400: the visual entries are duplicated under two naming conventions, so distinct modules are about 759, and half of 882 is 441. Sites: the comment above `IsFp8BlockProjection`, the comment above `Fp8BlockQuantConfig::modules_to_not_convert`, and `.agents/specs/model-fp8-block-weight.md`. The routing itself is correct and no defect in it is asserted; two other claims in the same comment were checked against the checkpoint headers and hold (zero `input_scale` tensors, and the `[96, 40]` block-grid hazard is real). Found while auditing the checkpoint for #1613, fixed in the same flow | bug | | [#1502](https://github.com/mudler/vllm.cpp/issues/1502) | `VT-REFTIER-HOST-ADDRESSABLE` | **`docs/ENVIRONMENT.md` described `VT_ADOPT_DEVICE_BYTES` as Vulkan-only and said it has "No effect on CUDA/CPU/Metal", and [`cffe59b02`](https://github.com/mudler/vllm.cpp/commit/cffe59b02) ([#1477](https://github.com/mudler/vllm.cpp/issues/1477)) made both halves false.** That change moved `ReferenceTierEligible` off `UnifiedMemory()` onto `Backend::DeviceMemoryIsHostAddressable()` and added truthful overrides so no backend lost the reference tier, so `MetalBackend` now answers `MetalContext::unified_memory()` and `RocmBackend` answers its `unified_memory_`. The weight loader gates the lever on exactly that predicate, at both `AdoptDeviceBytesAsHost` branches in `src/vllm/model_executor/models/qwen3_5_weights.cpp`, so the lever ACTS on Apple silicon and on an integrated ROCm part. **The correction is not "add two backend names".** Every number in that row is GB10 through Vulkan, and nobody has measured the lever on either new arm, so the row now separates the backends it is MEASURED on from the backends that merely satisfy the predicate — reach and measurement are different claims and the row read as if the measurement covered the reach. CUDA and CPU stay inert and are unchanged: neither overrides the default `false`, which `tests/vllm/platforms/test_platform.cpp` pins for GB10, and the CPU backend answering `UnifiedMemory() == true` while the narrower predicate stays `false` is the whole reason the two properties are separate. The MEASUREMENT on Metal and integrated ROCm stays owed and is listed under `## Owed` in [`vt-reference-tier-host-addressable.md`](specs/vt-reference-tier-host-addressable.md); it needs an Apple-silicon box or an integrated AMD part | documentation | +| [#1629](https://github.com/mudler/vllm.cpp/issues/1629) | `KERNEL-ATTN-DENSE-FLASH` | **`test_check_attention_rung_consistency.py` stored a count of the model tree, so every row on the attention-rung allowlist redded it by doing the thing the allowlist exists for.** `ShippedTreeTests::test_the_population_is_not_empty` asserted `>= 9` against a tree holding exactly 9 `vt::Attention(` sites, so a removing row had zero headroom and no green path: leaving the parked stem redded the floor (`8 not greater than or equal to 9`), and deleting it redded the floor and `test_allowlist_holds_only_the_in_flight_stems` as well -- while the allowlist header explicitly recommends the first of those two. That is the `## Records` shape AGENTS.md names, a measurement of one file stored inside another, and it blocked PR #1579 (#1545) and the LTX-2.5 routing row, which removes two of the three parked stems. FIXED HERE, and NOT by lowering the number, which is the known mute-switch: the floor became `>= 1`, because an empty population means the scanner broke and that is the only thing a raw total can honestly detect, and the guard that a rename cannot slip past stays `test_the_six_deliberate_sites_carry_a_marker`, which pins six sites BY NAME. A case that the stem in a red message names a real source file was added beside it, so a typo in the allowlist is still caught without pinning a count. A second drift lock in the same suite, `assertGreater(excused, 0)`, required the shipped allowlist to stay non-empty forever; it is replaced by two synthetic cases that build their own allowlisted population, so the excused counter is pinned without the shipped tree having to keep a stem parked. Found while landing #1578 and #1579 together -- each green in isolation, main red once both land -- and fixed in the same flow | bug | +| [#1631](https://github.com/mudler/vllm.cpp/issues/1631) | — | **A comment-only edit is impossible in any of the 43 `scripts/check-*.py` checkers, so a comment that is measurably false in one cannot be corrected.** `scripts/check-pr-size.py:170` classifies every `scripts/check-*.py` and `scripts/check-*.sh` as a `governance_checker`, and `change_errors` then demands a paired `tests/scripts/test_*.py` change that `executable_evidence` proves goes RED against the BASE checker. A comment-only diff leaves BASE and HEAD semantically identical, so no test can distinguish them and no such evidence can exist. Measured on this row: `ERROR: BASE checker stayed green for 'scripts/check-attention-rung-consistency.py'; changed test is not semantic evidence`, rc=1, with the identical invocation against the parent commit exiting 0. Live cost, three comments in `scripts/check-attention-rung-consistency.py` that ship unrepaired in #1578: `:58-61` says widening to `\bAttention\s*\(` is not the repair because it would match every fast rung, when the reason a wider pattern is not the repair is the function-pointer call it still cannot reach; `:93-96` says the `\b` is what excludes `vt::AttentionDenseFlash(`, when the trailing `\(` is, and the `\b` only excludes a leading identifier character as in `xyvt::Attention(`; `:252-255` says `sites - marked` is not the excused count, when on this tree it is (9 sites, 6 marked, 3 excused). The suite beside them was repaired for #1629, so the tree now contradicts itself across two files in the same directory pair. NOT fixed in the flow that filed it: teaching the guard to tell a comment-only or docstring-only diff from a semantic one changes what the gate accepts, which AGENTS.md `## Changing the rules or a checker` routes to its own row, spec and red-before evidence, and the honest report is therefore a filed gap rather than a comment smuggled in beside an unrelated semantic change. A candidate patch is parked on the issue, and two smaller pre-existing defects in `check-pr-size.py` itself (an incomplete entry-point list at `:370-371`, an unread `SELF_CHECKER` constant at `:376`, which the issue body records as `:378` because a line anchor drifts inside the pull request that writes it) are frozen by the same lock. Owed under `## Owed` in [attention-rung-visibility.md](specs/attention-rung-visibility.md) | bug | | [#1632](https://github.com/mudler/vllm.cpp/issues/1632) | `QUANT-QWEN38-27B-NVFP4-ARM` | **W6's NVFP4 token gate named [#1185](https://github.com/mudler/vllm.cpp/issues/1185) as the authority it waits on, and #1185 closed on 2026-08-18 as local-only** -- it tracked one operator's machines rather than a defect here -- so five sites pointed a reader at an issue that reports "closed" without reporting "cleared": `docs/FEATURES.md`, and the spec's `**Related:**` header, wave table, blockers section, `## Owed` list and `## Now`. **The blocker did not close with the issue, and it is not the one the citations described.** The pinned oracle `5559679229bc961848b121ccdeaa8fa5d79bec98` DOES build, install, import and GENERATE TOKENS inside an `rc` lease on `dgx:gpu0` (2026-08-18), which kills the "a model run is untested" clause those sites carried, and #1213 killed the "a lease cannot produce a runtime" premise underneath it. It survived at `max_num_batched_tokens` 512, `max_model_len` 512 and `gpu_memory_utilization` 0.30 on a ~20 GiB model, where the recorded denominator for this family is 8192 and 2048; `AGENTS.md` §Gates requires vLLM's PRODUCTION configuration as the denominator, so a reduced-`mnbt` arm is a different engine setup rather than a smaller measurement, and `gpu_memory_utilization` is a REFUTED lever (`.agents/specs/mtp-k-gt-1.md`: 0.75 thrashed 42 minutes, 0.30 rebooted the box). The named next levers are `max_num_batched_tokens` and `cudagraph_capture_sizes`, one at a time. The second half is the bytes: `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121`@`36f717a2` is ~20.4 GiB over four shards and is not mirrored where a lease can read it, which is also why its sha256 is recorded as unpaid. Same shape as [#1613](https://github.com/mudler/vllm.cpp/issues/1613) for the block-wise FP8 gate. FIXED IN FLOW: all five citations now name this issue, and the loader is untouched -- W5's accounting and cross-check need no lease and no oracle. Spec [`qwen38-27b-quant-arms.md`](specs/qwen38-27b-quant-arms.md), parent [#821](https://github.com/mudler/vllm.cpp/issues/821) | gap | | [#1538](https://github.com/mudler/vllm.cpp/issues/1538) | `SPEC-DFLASH2` | **vllm#52816's head moved a THIRD time (`66e5414c` -> `3406ec1d`) while it is still open, and refactors `compute_candidates` into `LogitsProcessor.get_top_k_tokens`.** Measured 2026-08-21 by W6 from `raw.githubusercontent.com` at both heads: +11/-80 on `qwen3_dflash2.py`, +4/-16 on `dflash2/speculator.py`, +2/-5 on the base `speculator.py`. The big one is a RELOCATION rather than new math -- the padding mask, the id rebase, the TP all-gather and the scale-THEN-softcap order all survive in `logits_processor.py:241-286`, so `## Owed` O16's reading of the codebook-span question holds at BOTH heads. What does NOT survive is the explicit `UnquantizedEmbeddingMethod`/`UnquantizedLinearMethod` guard that `## Risks/decisions` D12 ports as `RefuseQuantizedDflash2LmHead`, which is deleted at `3406ec1d`; our guard's own reason (the GGUF arm dequantizes `output.weight` to bf16, and a GGUF target with a safetensors DFlash2 draft is admitted here) is independent of upstream's and stands. NOT reconciled in flow, deliberately: `## Gates` G2 fixes the gate head at `66e5414c` while the pull request is unmerged, and moving the port onto a third unmerged head during the gate would move the thing being measured. Owed under `## Owed` O21 of [the DFlash2 spec](specs/dflash2-spec-decode.md) | verification | | [#1456](https://github.com/mudler/vllm.cpp/issues/1456) | `SPEC-DFLASH2` | **The GB10 oracle DOES have a FLASH_ATTN denominator: the arch measurement stands, the conclusion drawn from it does not.** Measured 2026-08-21 by W6 on `dgx:gpu0` through an `rc` lease, with the very wheel #1456 was filed about (`vllm-0.1.dev1+g66e5414c6`, sha256 `fbc247ab...`). A capture that exported `VLLM_ATTENTION_BACKEND=TRITON_ATTN` got `FLASH_ATTN` anyway and RAN: `Using FlashAttention version 2`, 54.87 GiB loaded, CUDA graphs captured (PIECEWISE 5/5, FULL 1/1, plus the DFlash2 speculator's own), 4 x 64 coherent tokens, speculation live at 209 accepted of 350 drafted and mean acceptance length 5.00. No `cudaErrorUnsupportedPtxVersion`. Consistent with the `sm_80`/`sm_75` SASS finding rather than contradicting it: `sm_80` PTX JITs FORWARD, and that error is the OPPOSITE failure (PTX newer than the driver). So `FA_USABLE=0` in the staged `FA-CONSTRAINT.txt` was inferred from emitted arches, never observed from a run, and is the thing to reconcile. A SECOND trap found in the same run and recorded so nobody repeats it: **`VLLM_ATTENTION_BACKEND` does not exist at this revision** -- grepping every `.py` in the wheel returns nothing; the knob is `EngineArgs.attention_backend` (`arg_utils.py:706`) folded into `AttentionConfig.backend` (`:2382`), so the old export selects NOTHING and auto-selection wins silently, letting a run record one backend while executing another. NOT reconciled in flow: W6 does not substitute a denominator the developer declared, and takes both arms instead, each named in its own golden. Owed under `## Owed` O22 of [the DFlash2 spec](specs/dflash2-spec-decode.md) | verification | diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 871b9cacf..193ada943 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -150,7 +150,7 @@ host/sched. Detail: state `KERNEL-FA2-GQA-SWAP-FLIP`. | `KERNEL-MOE-SQRTSOFTPLUS-HASH` | **DeepSeek-V4 MoE router + clamped-SwiGLU deltas — the three genuinely-new-vs-V2/V3 MoE pieces** (DeepSeek-V4-Flash W6). V4 keeps the DeepSeek grouped-GEMM / 256-expert w13/w2 / shared-expert / NVFP4 machinery (REUSED, not re-ported) but replaces three primitives. Three ops: **(1)** the router SCORE function **`sqrt(softplus(x))`** (`softplus(x)=log(1+exp(x))`, then sqrt) — distinct from V2/V3's sigmoid/softmax `noaux_tc`; the sqrt∘softplus COMPOSITION is load-bearing (RED-first proven); **(2)** the router: score all experts, add `e_score_correction_bias` for SELECTION ONLY, pick top-k OR — for the first `num_hash_layers` HASH layers — look experts up directly in the `tid2eid` [vocab, topk] token-id→expert table (BYPASSING top-k), GATHER weights from the UNBIASED scores, renormalize, ×`routed_scaling_factor` (the bias-affects-selection-not-weights split + the hash bypass are load-bearing, both RED-first proven); **(3)** the **clamped SwiGLU** expert activation `SiluAndMulWithClamp` — `gate` clamped `max=limit` (max ONLY), `up` clamped `[-limit,+limit]` (BOTH sides), then `gate·sigmoid(alpha·gate)·(up+beta)`; the ASYMMETRIC clamp is load-bearing (RED-first proven). MegaMoE (SM100-only) is NOT the GB10 target — this mirrors the FusedMoE-fallback router GB10 runs | score `vllm/model_executor/layers/fused_moe/router/fused_topk_bias_router.py:88`; router `fused_topk_bias_router.py:75-118` (`_topk_softplus_sqrt_torch`) + hash branch `:100-106` + dispatch `:254-265`; hash-table wiring `vllm/models/deepseek_v4/nvidia/model.py:562-578,:686,:696-717`; FusedMoE fallback `nvidia/model.py:647-691`; clamped SwiGLU `vllm/model_executor/layers/activation.py:197-201` (`SiluAndMulWithClamp.forward_native`), used by DeepseekV4MLP `nvidia/model.py:126-133`; cross-checked SGLang `v0.5.15` `python/sglang/srt/layers/moe/{topk.py:1013-1014, hash_topk.py:137-180}` @ `555967922` | Portable host reference (device kernels landed W7-device — see `KERNEL-DSV4-W7-DEVICE`) [deepseek_v4_moe.cpp](../src/vllm/model_executor/models/deepseek_v4_moe.cpp) + [deepseek_v4_moe.h](../include/vllm/model_executor/models/deepseek_v4_moe.h): `SqrtSoftplus` / `SqrtSoftplusRouteTopk` / `ClampedSwiGLU` | **CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, `-Wall -Werror -Wextra` 0-warn on the new TUs):** [test_deepseek_v4_moe.cpp](../tests/vllm/models/test_deepseek_v4_moe.cpp) **12/12 cases · 716 assertions** — hand-derived literal cases (sqrt∘softplus composition `softplus(x)=4 ⇒ score=2`; bias flips selection but weight stays the UNBIASED 1.0 not 3.0; renormalize by the unbiased sum; routed_scaling_factor; hash `tid2eid` picks {3,1} where top-k would pick {2,0}; asymmetric clamp gate=-5 kept vs up clamped to -2; gate/up clamp boundaries; alpha/beta) + from-first-principles double-precision references (router f32==f64 rel-L2 < 1e-5 + exact ids; SqrtSoftplus f64 + monotonicity; ClampedSwiGLU rel-L2 < 1e-6). **RED-first PROVEN all three levers:** drop the sqrt → 8 cases/493 assertions fail; gather weights from the BIASED scores → 2 cases/181 fail; symmetric-clamp the gate → 2 cases/6 fail; revert restores 12/12·716. Honest gate form: host-reference + hand-case + structural review vs vLLM+SGLang `file:line` (fixed-config 167B not constructible at a tiny shape ⇒ NOT a dumped-oracle rel-L2). Full-model gate multi-Spark-blocked (156.7 GiB); the device kernels reuse the existing grouped-GEMM + `DeepseekV4Model::Forward` assembly (W7) + the strict/near-tie engine gate (W8) are named residuals | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W6 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W6` | | `KERNEL-DSV4-W7-DEVICE` | **DeepSeek-V4-Flash W7-DEVICE — the four NEW V4 op families' CUDA kernels**, each a 1:1 DEVICE port of the landed portable HOST reference (the oracle the four SPIKE rows above pinned), registered through the vt **OpProvider seam** (`kDeepseekV4{Mhc,Dsa,Compressor,Moe}`) so `DeepseekV4Model::ForwardDevice` can dispatch them: **(MHC)** Sinkhorn + mHC pre/post + hc_head collapse; **(DSA)** indexer weight-fold + weighted-MQA ReLU logits + causal top-k + per-head attention-sink softmax + grouped output-LoRA; **(Compressor)** softmax-window pool + RMSNorm + save-time APE + **fp8_ds_mla** KV encode (UE8M0 block scale + e4m3, bf16 rope) / decode; **(MoE)** sqrtsoftplus/hash router + clamped SwiGLU. The 512-wide MLA attention + expert grouped-GEMM REUSE the existing NVFP4/FP8 kernels (`kMlaDecodeAttention`/`kMoeGroupedGemmNvfp4`, `cuda_mla_attn.cu`/`cuda_moe*.cu`) and are NOT re-ported | the SAME `file:line` the host refs cite (the `KERNEL-{MHC-SINKHORN,ATTN-DSA-SPARSE-INDEX,ATTN-DSA-COMPRESSOR,MOE-SQRTSOFTPLUS-HASH}` upstream columns), @ `555967922` | [cuda_deepseek_v4.cu](../src/vt/cuda/cuda_deepseek_v4.cu) (kernels + host-vector launchers + OpProvider registration) + [deepseek_v4_device.h](../include/vllm/model_executor/models/deepseek_v4_device.h) / [deepseek_v4_device.cpp](../src/vllm/model_executor/models/deepseek_v4_device.cpp) (seam resolvers); `DeepseekV4Model::ForwardDevice` composes them ([deepseek_v4.cpp](../src/vllm/model_executor/models/deepseek_v4.cpp)) | **DGX GB10 (sm_121a) UNIT GATE GREEN + RUNTIME-VERIFIED (2026-07-29):** [test_cuda_deepseek_v4.cpp](../tests/vllm/models/test_cuda_deepseek_v4.cpp) **11/11 cases · 153 assertions** — each device kernel vs its host-ref oracle at small shape: BIT-EXACT ids (DSA causal top-k, sqrtsoftplus/hash router selection), `-inf` mask exact (indexer out-of-window), near-tie rel-L2 < 1e-4 for the fp reductions (Sinkhorn, pool/softmax, sqrtsoftplus — device `expf`/`sqrtf`/`rsqrt` vs host), fp8_ds_mla encode→decode within the e4m3 granularity bound, bf16 rope bit-exact; PLUS the **ForwardDevice composition gate** (device forward == host forward, rel-L2 < 2e-3 over the 4-family tiny-config interleave). **compute-sanitizer memcheck 0 errors.** **RED-first PROVEN:** dropping the sqrt in the device sqrtsoftplus fails 3 cases / 6 assertions (sqrtsoftplus + router weights + ForwardDevice); revert restores 11/11·153. Build: CUDA `-Werror` clean (the #155 voxtral GCC-13 `-O2` array-bounds/stringop false positive neutralized locally). Honest 3-state: RUNTIME-VERIFIED at small shape on real GB10; the real-checkpoint paged-engine e2e stays W8 (156.7 GiB does not fit ONE GB10). **DECODE GLUE-FOLD LANDED (2026-08-03, GB10 sm_121a):** the resident-decode `norm_rope_rows` + MHC-pre kernels folded FP64→FP32 — `VT_V4_ROPE_FLOAT` (fused norm+RoPE **4.58→0.46 ms/step ~10×**, decode +6.1%) + `VT_V4_MHC_LEAN` (finish block 256→1024 + sqrsum-fold, +0.7%, floored by 86 sequential single-block launches/step), both default-ON + BYTE-EXACT (decode ids `=1`/`=0` token-identical via the resident-decode path `--gpu --kv-cache`); `test_cuda_deepseek_v4` Brick-7 + Brick-B **20/20·67073** PASS on GB10; net decode 14.02→14.96 tok/s → 90.7% of ds4 ~16.5 (`CLAIM-DSV4-ROPE-FLOAT` / `CLAIM-DSV4-MHC-LEAN`; ds4 bar corrected from the unreproduced 17.13 anchor to the fair same-session ~16.5, the later MHC-SINK4 reached ~96%) | [deepseek-v4-flash spike](specs/deepseek-v4-flash.md) §W7 | `SPIKE` | `CLAIM-DEEPSEEK-V4-W7-DEVICE` | | `KERNEL-KDA-DELTA` | **Kimi Delta Attention (KDA) gated-linear-attention delta vs plain GDN — a genuinely new gated-linear-attention family, the shared unblocker for Kimi-Linear-48B and Kimi-K3 (W4).** `KimiGatedDeltaNetAttention` SUBCLASSES `GatedDeltaNetAttention`, so its conv-state/cache layout, `GDNAttentionMetadata`, chunked-delta recurrence and WY solve are REUSED from our landed GDN — this row owns ONLY the four KDA-specific deltas plain GDN lacks: **(1)** a per-channel **`[H,D]` low-rank decay** via an `f_a_proj→f_b_proj` bottleneck (GDN has only a per-HEAD scalar decay from `A_log`); **(2)** the decay GATE `g = -exp(A_log[h])·softplus_β(g1+dt_bias)` per channel (β=1, thr=20; `kda_gate_fwd_kernel` decode) + its chunk-local cumulative-sum prefill variant (`kda_gate_cumsum_fwd_kernel`, folds `RCP_LN2`); **(3)** the **sigmoid-gated output norm** `FusedRMSNormGated(head_dim, activation="sigmoid")` = `rmsnorm(x)·w·σ(g)` (the gated-linear-attention output norm GDN lacks); **(4)** three separate q/k/v short causal convs (`conv_size=4`, silu) + the q/k **L2-norm** preprocessing (`x/sqrt(Σx²+eps)`, SUM not mean). ADDITIVE — does NOT touch `cuda_gdn.cu`/`gdn_attn.cpp`, so the Qwen3.6-27B/35B GDN gate is structurally untouched (like DSA kept shared-MLA untouched) | decay bottleneck `vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py:142-156,:245`; decay gate `vllm/third_party/flash_linear_attention/ops/kda.py:1541-1600,:1603-1646`; chunk-cumsum `kda.py:1182-1254,:1257-1303`; gated norm `kda.py:463-487` (`:436` eps=1e-5); short conv `kimi_gdn_linear_attn.py:171-198,:324-356`; q/k L2-norm `kda.py:1511-1513` + `ops/l2norm.py:42-43,:96` @ `555967922` | Portable host reference (device kernel is a named residual) [kimi_kda.cpp](../src/vllm/model_executor/models/kimi_kda.cpp) + [kimi_kda.h](../include/vllm/model_executor/models/kimi_kda.h): `KdaLowRankDecay` / `KdaDecayGate` / `KdaDecayGateChunkCumsum` / `FusedRMSNormGated` / `KdaShortConv` / `L2NormRows` | **CPU UNIT GATE GREEN (2026-07-28, `-Wall -Werror -Wextra` 0-warn):** [test_kimi_kda.cpp](../tests/vllm/models/test_kimi_kda.cpp) **14/14 cases · 36 assertions** — hand-derived literal cases (f_b∘f_a bottleneck; `-exp(A_log)·softplus` with the >thr linearisation; per-head A_log + per-channel dt_bias; chunk-cumsum reset+`RCP_LN2` fold; sigmoid-gated norm; swish-vs-sigmoid branch; per-head-dim normalisation; causal-depthwise+silu conv; zero-init-state edge; L2-norm SUM-not-mean) + from-first-principles double-precision references on randomized shapes (decay gate, gated norm, short conv rel-L2 < 1e-6). Honest gate form: host-reference + structural review, NOT a dumped-oracle rel-L2 — the REAL e2e gate is the Kimi-Linear-48B-A3B proxy vs the pinned oracle (DGX-blocked; K3 2.8T does not fit one GB10). Named residuals: the KDA CUDA device kernel + the Kimi-Linear-48B proxy gate — anchor `tests/vllm/models/test_kimi_kda.cpp:41` | [kda-kernel-delta spike](specs/kda-kernel-delta.md) | `SPIKE` | `CLAIM-KDA-KERNEL` | -| `KERNEL-ATTN-DENSE-FLASH` | **Flash-TILED dense non-causal attention — the SHARED-MEMORY-TILED form of `AttentionDenseFast` for long non-causal contexts** (multimodal-speed §14, the Whisper AUDIO encoder — hd-64, non-causal, 1500 frames × 32 layers). A block of `kFlashBr=16` query-warps (512 threads) SHARES each streamed `kFlashBc=64`-column K/V tile out of shared memory (classic FlashAttention K/V tiling): the CTA cooperatively loads a K/V tile into shared memory, then each warp runs its online-softmax update reading K/V from shared memory, killing `AttentionWarpKernel`'s O(t²) redundant global K/V re-reads (one full K/V sweep per (query,head)). One q-head per CTA (all warps share the GQA kv-head). BIT-IDENTICAL to `AttentionDenseFast`: the per-warp arithmetic (per-lane head_dim grouping `lane+32k`, butterfly `__shfl_xor`, sequential j-order, f32 online-softmax `m`/`l`/`acc`) is copied verbatim, only K/V bytes come from shared memory instead of global ⇒ token-identical by construction. **Head_dim-generic** (`npl=(d+31)/32`, d≤256): since 2026-07-28 (multimodal-speed §16, `CLAIM-MM-SPEED-QWEN-IMAGE`) ALSO the default for the Qwen3-VL / Qwen3.6-27B VISION tower per-frame self-attention (hd-72, non-causal, 784 patches) — byte-identical to the warp `AttentionDenseFast` it replaced there (bench 0/1,003,520 mismatch; STRICT image/video e2e 32/32) | STRUCTURE ported 1:1 from vendored FlashAttention-2 `compute_attn_1rowblock` [flash_fwd_kernel.h:52](../src/vt/cuda/flash_attn/src/flash_fwd_kernel.h#L52) (sK/sV shared tiles :163-165 + the `for(int n_block…)` K/V-tile stream + online rescale); non-causal encoder dispatch cross-checked to vLLM `WhisperEncoderAttention` [whisper.py:255](https://github.com/vllm-project/vllm/blob/e24d1b24/vllm/model_executor/models/whisper.py#L255) | `OpId::kAttentionDenseFlash` + decl [ops.h](../include/vt/ops.h) + wrapper/validation [ops.cpp](../src/vt/ops.cpp); CUDA `AttentionDenseFlashKernel`/`AttentionDenseFlashKernelCuda` [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu); CPU maps to `AttentionKernel` (byte-identical) [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp); wired default in [whisper_audio.cpp](../src/vllm/model_executor/models/whisper_audio.cpp) (`VT_WHISPER_ENC_WARP=1`/`VT_WHISPER_ENC_EAGER=1` A/B) + [qwen3_vl_vision.cpp](../src/vllm/model_executor/models/qwen3_vl_vision.cpp) (vision tower default since §16; `VT_QWEN3VL_ATTN_WARP=1`/`VT_QWEN3VL_ATTN_EAGER=1` A/B) | **GPU GATE GREEN on dgx (2026-07-28, GB10 sm_121a, base `af1ed76b`):** CUDA `-Werror` 0-warn (cutlass-ON + FA2-ENABLED banners). `test_voxtral_e2e` **16/16** default-flash; flash/warp/eager token dumps md5-IDENTICAL (`89923566…`) ⇒ ZERO token flips; goldens md5 UNCHANGED (`voxtral_golden.json 8ab87b7e…`, `voxtral_neartie.json 937b9ad3…`, before==after). Proof-of-run nsys `AttentionDenseFlashKernel` 32 inst, ZERO `AttentionWarpKernel`/naive on encoder; RED confirmed (corrupt kernel → gate FAILS → restore → 16/16); `compute-sanitizer --tool memcheck` **0 errors**; 3 runs byte-identical. **A/B (same binary, `flock`, rep0 dropped):** attention **35.11 → 19.29 ms/layer (1.82×, NON-OVERLAPPING)**; encoder forward **~1834 → ~1375 ms (1.33×)**. **NOT at parity:** ~1.37 s vs vLLM ~43 ms TTFT (~32×, was ~44×) — the scalar warp-per-query recurrence is now serial-latency-bound over 1500 keys (L2 already served much of the redundant reads ⇒ 1.8× not 16×); gap-closer is a tensor-core MMA hd-64 non-causal FA2 instantiation (LARGE) + resident encoder weights (MEDIUM). **Vision tower (§16, 2026-07-28):** extended to the Qwen3-VL/27B tower (hd-72, 784 patches) — STRICT image/video e2e 32/32, bench flash-vs-warp 0/1,003,520 mismatch, nsys default 4B e2e `AttentionDenseFlashKernel` 24 inst/zero warp, RED 30/46→46/46, sanitizer 0; A/B warp 148.3→flash 142.3 ms = 1.04× (small — the vision attention at t=784 is serial-latency-bound not bandwidth-bound; the tower already BEATS vLLM at 0.57× eager) | [multimodal-speed](specs/multimodal-speed.md) §14 + §16 | `ACTIVE` | `CLAIM-MM-SPEED-AUDIO-ENC-KERNEL` + `CLAIM-MM-SPEED-QWEN-IMAGE` | +| `KERNEL-ATTN-DENSE-FLASH` | **Flash-TILED dense non-causal attention — the SHARED-MEMORY-TILED form of `AttentionDenseFast` for long non-causal contexts** (multimodal-speed §14, the Whisper AUDIO encoder — hd-64, non-causal, 1500 frames × 32 layers). A block of `kFlashBr=16` query-warps (512 threads) SHARES each streamed `kFlashBc=64`-column K/V tile out of shared memory (classic FlashAttention K/V tiling): the CTA cooperatively loads a K/V tile into shared memory, then each warp runs its online-softmax update reading K/V from shared memory, killing `AttentionWarpKernel`'s O(t²) redundant global K/V re-reads (one full K/V sweep per (query,head)). One q-head per CTA (all warps share the GQA kv-head). BIT-IDENTICAL to `AttentionDenseFast`: the per-warp arithmetic (per-lane head_dim grouping `lane+32k`, butterfly `__shfl_xor`, sequential j-order, f32 online-softmax `m`/`l`/`acc`) is copied verbatim, only K/V bytes come from shared memory instead of global ⇒ token-identical by construction. **Head_dim-generic** (`npl=(d+31)/32`; the register blocking allows d≤256 but the K/V tile's dynamic shared memory is what BINDS — d≤192 bf16 / d≤96 f32 under CUDA's default 48 KiB cap, see the 2026-08-21 entry): since 2026-07-28 (multimodal-speed §16, `CLAIM-MM-SPEED-QWEN-IMAGE`) ALSO the default for the Qwen3-VL / Qwen3.6-27B VISION tower per-frame self-attention (hd-72, non-causal, 784 patches) — byte-identical to the warp `AttentionDenseFast` it replaced there (bench 0/1,003,520 mismatch; STRICT image/video e2e 32/32) | STRUCTURE ported 1:1 from vendored FlashAttention-2 `compute_attn_1rowblock` [flash_fwd_kernel.h:52](../src/vt/cuda/flash_attn/src/flash_fwd_kernel.h#L52) (sK/sV shared tiles :163-165 + the `for(int n_block…)` K/V-tile stream + online rescale); non-causal encoder dispatch cross-checked to vLLM `WhisperEncoderAttention` [whisper.py:255](https://github.com/vllm-project/vllm/blob/e24d1b24/vllm/model_executor/models/whisper.py#L255) | `OpId::kAttentionDenseFlash` + decl [ops.h](../include/vt/ops.h) + wrapper/validation [ops.cpp](../src/vt/ops.cpp); CUDA `AttentionDenseFlashKernel`/`AttentionDenseFlashKernelCuda` [cuda_ops.cu](../src/vt/cuda/cuda_ops.cu); CPU maps to `AttentionKernel` (byte-identical) [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp); wired default in [whisper_audio.cpp](../src/vllm/model_executor/models/whisper_audio.cpp) (`VT_WHISPER_ENC_WARP=1`/`VT_WHISPER_ENC_EAGER=1` A/B) + [qwen3_vl_vision.cpp](../src/vllm/model_executor/models/qwen3_vl_vision.cpp) (vision tower default since §16; `VT_QWEN3VL_ATTN_WARP=1`/`VT_QWEN3VL_ATTN_EAGER=1` A/B) | **GPU GATE GREEN on dgx (2026-07-28, GB10 sm_121a, base `af1ed76b`):** CUDA `-Werror` 0-warn (cutlass-ON + FA2-ENABLED banners). `test_voxtral_e2e` **16/16** default-flash; flash/warp/eager token dumps md5-IDENTICAL (`89923566…`) ⇒ ZERO token flips; goldens md5 UNCHANGED (`voxtral_golden.json 8ab87b7e…`, `voxtral_neartie.json 937b9ad3…`, before==after). Proof-of-run nsys `AttentionDenseFlashKernel` 32 inst, ZERO `AttentionWarpKernel`/naive on encoder; RED confirmed (corrupt kernel → gate FAILS → restore → 16/16); `compute-sanitizer --tool memcheck` **0 errors**; 3 runs byte-identical. **A/B (same binary, `flock`, rep0 dropped):** attention **35.11 → 19.29 ms/layer (1.82×, NON-OVERLAPPING)**; encoder forward **~1834 → ~1375 ms (1.33×)**. **NOT at parity:** ~1.37 s vs vLLM ~43 ms TTFT (~32×, was ~44×) — the scalar warp-per-query recurrence is now serial-latency-bound over 1500 keys (L2 already served much of the redundant reads ⇒ 1.8× not 16×); gap-closer is a tensor-core MMA hd-64 non-causal FA2 instantiation (LARGE) + resident encoder weights (MEDIUM). **Vision tower (§16, 2026-07-28):** extended to the Qwen3-VL/27B tower (hd-72, 784 patches) — STRICT image/video e2e 32/32, bench flash-vs-warp 0/1,003,520 mismatch, nsys default 4B e2e `AttentionDenseFlashKernel` 24 inst/zero warp, RED 30/46→46/46, sanitizer 0; A/B warp 148.3→flash 142.3 ms = 1.04× (small — the vision attention at t=784 is serial-latency-bound not bandwidth-bound; the tower already BEATS vLLM at 0.57× eager) **2026-08-21 (`CLAIM-ATTN-RUNG-VISIBLE`, issue [#1544](https://github.com/mudler/vllm.cpp/issues/1544), spec [attention-rung-visibility.md](specs/attention-rung-visibility.md)): the advertised head_dim contract is now the LAUNCHABLE one, and the naive rung stops being a silent default.** The op stated `d <= 256` ([cuda_ops.cu](../src/vt/cuda/cuda_ops.cu) `LaunchAttentionDenseFlash`) while requesting `2*kFlashBc*d*sizeof(Tin)` bytes of DYNAMIC shared memory with no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`, so the driver's default 48 KiB cap made the real ceiling **192 bf16 / 96 f32** — Kimi (192 f32, 96 KB) and Qwen3.5 (256) would have taken a bare launch error from the `cudaGetLastError` at the bottom of the launcher, naming nothing they could do instead. The bound now lives in [ops.h](../include/vt/ops.h) as `AttentionDenseFlashSmemBytes` / `AttentionDenseFlashMaxHeadDim`, PURE host arithmetic so a box with no GPU can execute it, tied to the kernel by two `static_assert`s on `kFlashBc` and the register blocking; the launcher refuses above it naming `vt::AttentionDenseFast`, which uses NO shared memory and does serve those widths. NARROWING was chosen over `cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize)`: d=256 f32 wants 128 KiB, above the opt-in per-block cap of the consumer Blackwell parts gated here, so the opt-in would still leave the widest advertised width a lie AND cannot be verified without a device. Strictly additive for callers — the bound is INCLUSIVE, so d=192 bf16 lands exactly on 49152 and still launches. Mirrors `supports_head_size` / `get_supported_head_sizes` [backend.py:155-163](https://github.com/vllm-project/vllm/blob/555967922/vllm/v1/attention/backend.py#L155), consulted BEFORE dispatch rather than discovered by launching. Same change adds [check-attention-rung-consistency.py](../scripts/check-attention-rung-consistency.py) (preflight + CI): a model TU naming `vt::Attention` needs a `// 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. `kAttention` and every existing caller's numerics are UNTOUCHED by construction — the checker executes no model code and the head_dim guard only fires where the launch already failed. CPU-GATED: checker green (9 sites / 6 marked / 3 unmarked and excused by the 3 in-flight allowlisted stems), unit + mutation cases in [test_check_attention_rung_consistency.py](../tests/scripts/test_check_attention_rung_consistency.py) (no case count is recorded here -- a count of one file stored in another is a drift lock; run the suite for the live number), new head_dim contract cases in [test_ops_attention.cpp](../tests/vt/test_ops_attention.cpp). OWED [#1573](https://github.com/mudler/vllm.cpp/issues/1573): the on-device refusal case and its reachability mutation are PENDING a lease — `dgx:gpu0` was held by the developer, and the CPU cases pin the arithmetic, never that the launcher calls it. | [multimodal-speed](specs/multimodal-speed.md) §14 + §16 | `ACTIVE` | `CLAIM-MM-SPEED-AUDIO-ENC-KERNEL` + `CLAIM-MM-SPEED-QWEN-IMAGE` + `CLAIM-ATTN-RUNG-VISIBLE` | | `KERNEL-MOE-ROUTING` | Router top-k, align, permute/unpermute, combine, activation | core MoE sources `CMakeLists.txt:1135-1157`; M=1 decode parallelization mirrors `topk_softmax_kernels.cu:192-242,494-537` (moeTopK/topkGating) + `moe_align_sum_kernels.cu:147-185,295-324`; **grouped-topk (`noaux_tc`)** `fused_moe/router/grouped_topk_router.py:106-161` (`forward_native`; the fused `ops.grouped_topk` at `:28-70` is the same formula), upstream tests `tests/kernels/moe/test_grouped_topk.py`, `test_routing.py` | [cuda_moe.cu:349](../src/vt/cuda/cuda_moe.cu#L349); parallel router argmax [cuda_moe.cu:61](../src/vt/cuda/cuda_moe.cu#L61); parallel moe_align BlockScan [cuda_marlin_repack.cu:224](../src/vt/cuda/cuda_marlin_repack.cu#L224); **grouped-topk (MLA campaign W3)** — additive `MoeRouterTopKArgs` fields + optional `e_score_correction_bias` arg [ops.h](../include/vt/ops.h), CPU ref `MoeRouterGroupedTopKKernel` [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp) + CUDA `MoeRouterGroupedTopKKernel` [cuda_moe.cu](../src/vt/cuda/cuda_moe.cu). A SEPARATE kernel: `num_expert_group == 0` still dispatches the original, so the existing router is byte-identical by construction | [routing tests](../tests/vt/test_ops_moe.cpp#L57); byte-exact router+align parity [test_ops_moe_grouped.cpp:451](../tests/vt/test_ops_moe_grouped.cpp#L451); 35B 315/315 gate; **grouped-topk** [test_ops_moe_router_grouped.cpp](../tests/vt/test_ops_moe_router_grouped.cpp) at REAL DeepSeek-V3 dims (256 experts, n_group=8, topk_group=4, top_k=8, sigmoid, routed_scaling 2.5, WITH `e_score_correction_bias`) vs an INDEPENDENT sort-based transcription of the upstream formula, plus isolated cases for bias-selects/unbiased-weights, top-2-sum-vs-max group scoring, the group mask excluding the global argmax, and renorm-before-scaling; CPU-vs-CUDA ids EXACT + run-to-run bit-reproducible | [inventory](specs/kernel-family-inventory.md) | `ANCHOR-BACKFILL` | `CLAIM-MOE-DECODE-PARALLEL-1` | | `KERNEL-MOE-UNQUANTIZED` | Unquantized grouped/batched MoE GEMM | core MoE sources `CMakeLists.txt:1135-1157`; upstream `tests/kernels/moe/test_unquantized_backend_selection.py` | activation/combine subset [cuda_moe.cu:349](../src/vt/cuda/cuda_moe.cu#L349); grouped execution remains NVFP4-specialized | [MoE tests](../tests/vt/test_ops_moe.cpp#L193), [grouped tests](../tests/vt/test_ops_moe_grouped.cpp#L160) | [inventory](specs/kernel-family-inventory.md) | `PARTIAL` | - | | `KERNEL-MOE-QUANTIZED` | FP8/INT8/NVFP4/MXFP4 grouped MoE | CUTLASS/FP4 builds `CMakeLists.txt:865-1002`; NVFP4 oracle `fused_moe/oracle/nvfp4.py:38-276` | NVFP4 fallback [cuda_matmul_nvfp4.cu:761](../src/vt/cuda/cuda_matmul_nvfp4.cu#L761), Marlin [cuda_moe_marlin.cu:156](../src/vt/cuda/cuda_moe_marlin.cu#L156) | [NVFP4 grouped tests](../tests/vt/test_ops_moe_grouped.cpp#L160); 35B gate | [inventory](specs/kernel-family-inventory.md) | `PARTIAL` | - | diff --git a/.agents/specs/attention-rung-visibility.md b/.agents/specs/attention-rung-visibility.md new file mode 100644 index 000000000..1a7a94afd --- /dev/null +++ b/.agents/specs/attention-rung-visibility.md @@ -0,0 +1,322 @@ +# Attention rung visibility: the naive kernel stops being the silent default + +Issue: [#1544](https://github.com/mudler/vllm.cpp/issues/1544). +Owning row: KERNEL-ATTN-DENSE-FLASH (kernel-matrix.md), the row that already owns +`AttentionDenseFlash` and its head-dim contract. This spec is an increment on that +row, linked from its evidence cell, exactly as `fusion-consistency-audit.md` is an +increment on the fusion framework row. It is deliberately NOT added to that row's +`Spec` column, because `check-gate-commands.py` classifies a row from the FIRST +spec link in that column and the row is pinned there as `no-gates-section`; +promoting a second spec into that column would move the row into +`RUNNABLE_BASELINE` as a side effect of an unrelated change. + +Note on spelling: this document writes the owning row id WITHOUT backticks for the +same reason. `check-agent-record.py::check_spec` selects a row's governing spec by +searching for the backticked token, and a second spec carrying it can change which +file is held to the structured-section contract. + +## Scope + +Two additive changes, both from #1544's `## Owed`. Neither changes what any +existing caller computes. + +IN SCOPE: + +1. A checker, `scripts/check-attention-rung-consistency.py`, that refuses a model + translation unit which names `vt::Attention` — the naive, correctness-grade + rung — without a recorded reason beside the call. +2. The marker comments the six deliberate call sites already deserve, in + `whisper_audio.cpp`, `qwen3_vl_vision.cpp`, `kimi_linear_device.cpp`, + `qwen3_5.cpp`, `nemotron_h.cpp` and `nemotron_h_device.cpp`. +3. A repair of `AttentionDenseFlash`'s advertised head-dim contract in + `src/vt/cuda/cuda_ops.cu`, so the bound it states is the bound it can launch. +4. The shared-memory bound as a pure, unit-testable host function in + `include/vt/ops.h`, so the arithmetic is executable on a box with no GPU. + +OUT OF SCOPE, and each for a stated reason: + +- `vt::Attention`'s behaviour and `OpId::kAttention`'s resolution. Both are frozen + so text decode stays byte-identical (`src/vt/cuda/cuda_ops.cu:3120-3122`), and + six model sites use the naive kernel as the REFERENCE arm of a numeric gate or + an A/B knob. Auto-routing it would delete those reference arms. +- `ltx2.cpp`, `ltx2_device.cpp` and `muse_glimmer_vision.cpp`. Their `vt::Attention` + calls are being REMOVED by other rows in flight ([#1545](https://github.com/mudler/vllm.cpp/issues/1545) + for Muse Glimmer, the LTX-2.5 routing row for the other two). Editing the lines + those changes replace would conflict for no gain, so the three stems are carried + on `scripts/attention-rung-allowlist.txt` with the owning issue named. +- Opting `AttentionDenseFlash` in to the >48 KiB shared-memory cap. See + `## Risks/decisions` D2. + +## Upstream chain + +vLLM never lets a model reach a kernel by omission, and it never advertises a +head-size domain it cannot serve. Both halves of this row mirror that polarity at +the pinned oracle `555967922`. + +| Concern | vLLM at the pin | Ours before this row | +|---|---|---| +| Who picks the encoder attention kernel | `vllm/model_executor/models/vision.py:99` `get_vit_attn_backend(head_size, dtype)` — every vision tower ASKS a selector, which reads the shape and the override | the C++ function name the author typed; `vt::Attention` resolves straight to the naive kernel and nothing routes it up | +| What a backend's head-size domain means | `vllm/v1/attention/backend.py:155-163` — a backend DECLARES `get_supported_head_sizes()` and `supports_head_size()` is consulted BEFORE dispatch | `AttentionDenseFlash` declares `head_dim <= 256` and cannot launch above 192 (bf16) / 96 (f32) | + +We cannot mirror `get_vit_attn_backend` directly: our seam has no config object and +no per-model backend enum, and the deliberate frozen-reference arms have no vLLM +counterpart. What transfers is the PROPERTY the selector gives upstream for free — +the kernel a model runs is a declared choice, not an omission. A checker is how a +C++ tree without a selector gets that property, and item 3 is simply the +`supports_head_size` half stated truthfully. + +## Our baseline + +Measured on this tree at `04f1cead6`. + +Nine live `vt::Attention(` call sites under `src/vllm/model_executor/models/`: + +| File:line | Disposition | +|---|---| +| `whisper_audio.cpp:324` | A/B rung, `VT_WHISPER_ENC_EAGER=1`, default is `AttentionDenseFlash` | +| `qwen3_vl_vision.cpp:527` | A/B rung, `VT_QWEN3VL_ATTN_EAGER=1`, default is `AttentionDenseFlash` | +| `kimi_linear_device.cpp:598` | behind `VT_KIMI_DEVICE_MLA`, default off, recorded as a measured negative | +| `qwen3_5.cpp:5279` | reference arm; non-test callers are in `tests/` | +| `nemotron_h.cpp:671` | reference arm | +| `nemotron_h_device.cpp:330` | reference arm | +| `muse_glimmer_vision.cpp:639` | THE defect, #1545, another row in flight | +| `ltx2.cpp:959`, `ltx2_device.cpp:421` | THE defect, LTX-2.5 routing row in flight | + +`AttentionDenseFlash` at `src/vt/cuda/cuda_ops.cu:3336` states `d <= 256` and +requests `2 * kFlashBc * d * sizeof(Tin)` bytes of dynamic shared memory at `:3338` +with `kFlashBc = 64`. `grep -rn cudaFuncSetAttribute src/vt/cuda/` returns nothing, +so the launch is bounded by CUDA's default 48 KiB dynamic cap: + +| Input dtype | Bytes per head_dim | Largest head_dim that launches | Advertised | +|---|---:|---:|---:| +| bf16 | 256 | 192 | 256 | +| f32 | 512 | 96 | 256 | + +It fails LOUD — `Check(cudaGetLastError(), "attention-dense-flash launch")` at +`:3352` throws — so this is a wrong contract and a trap, never silent corruption. + +## Port map + +| Change | Path | +|---|---| +| Head-dim bound as pure host arithmetic | `include/vt/ops.h`, beside the `AttentionDenseFlash` declaration | +| Honest refusal at the launcher | `src/vt/cuda/cuda_ops.cu` `LaunchAttentionDenseFlash` | +| Rung-visibility checker | `scripts/check-attention-rung-consistency.py` | +| One source of truth for the tile bytes | `LaunchAttentionDenseFlash` now sizes its `shmem` request from the SAME `AttentionDenseFlashSmemBytes` the guard reads, so the two cannot disagree | +| In-flight stems, with owning issue | `scripts/attention-rung-allowlist.txt` | +| Marker comments | the six deliberate model translation units | +| Gate wiring | `scripts/agent-preflight.sh`, `.github/workflows/ci.yml` | + +The checker reuses `scripts/checker_text.py::normalize_source`, so a +commented-out, `#if 0`-ed or `if (false)`-ed call is a deletion to it and never a +site, and the reported `file:line` still describes the original file. It scans +both `*.cpp` and `*.h` under the two model directories, because a call moved into +an inline function or a template in a header would otherwise leave it green, and +it keys its results on the PATH rather than the file stem, because `ltx2.cpp` and +`ltx2.h` share a stem and one would overwrite the other. The allowlist still +matches on the stem, so one entry covers a model's whole translation unit. + +## Tests to port + +There is nothing to port. vLLM's `supports_head_size` contract is enforced by its +selector rather than by a test that pins the arithmetic, and no upstream test +covers a shared-memory ceiling that only exists in our scalar kernel. + +New, all runnable with no GPU: + +| Test | Pins | +|---|---| +| `tests/scripts/test_check_attention_rung_consistency.py` | the checker's pure functions, and six mutations that must go RED | +| `tests/vt/test_ops_attention.cpp` new cases | the shared-memory arithmetic, both honest bounds, and that 256 is outside both | + +One test needs a device and is declared PENDING rather than skipped quietly: a +CUDA case that calls `vt::AttentionDenseFlash` at head_dim 256 and requires the +refusal to name `AttentionDenseFast`. It emits a loud MESSAGE and returns on a box +with no CUDA backend, which is this box. + +## Gates + +Every leg below is CPU-only except where it says otherwise. There is no GPU lease +on this row: `dgx:gpu0` is held by the developer, and AGENTS.md forbids reaching a +fleet device outside a lease. + +1. **Ran:** `python3 scripts/check-attention-rung-consistency.py` reports zero + drift on the tree. +2. **Ran:** `python3 tests/scripts/test_check_attention_rung_consistency.py`. +3. **Ran:** `ctest -R test_ops_attention`, including the new head-dim contract + cases. +4. **Ran:** `scripts/agent-preflight.sh --staged`. +5. **Owed:** the on-device refusal case. It needs a CUDA backend, and it is listed + under `## Owed` with the issue that carries it. + +## Dependencies + +None on other rows. The three allowlisted stems depend on rows in flight only in +the sense that removing their entries is those rows' cleanup, and a stale entry is +reported and never a failure. + +## Work breakdown + +1. Spec, committed first. +2. The checker, its allowlist and its mutation suite. +3. Marker comments at the six deliberate sites. +4. The head-dim bound helper, the launcher refusal, and the contract tests. +5. Record edits: the owning row's evidence cell, the issue index, the claim. + +## Risks/decisions + +**D1 — a checker, not a selector.** #1544 owes "a selector or a checker" and both +are legitimate. A selector that auto-routes `vt::Attention` by shape was REJECTED, +and not on taste: six of the nine sites exist precisely BECAUSE they are the naive +kernel. `nemotron_h.cpp:671`, `nemotron_h_device.cpp:330` and `qwen3_5.cpp:5279` +are reference arms that tests compare against, and `whisper_audio.cpp:324` and +`qwen3_vl_vision.cpp:527` are the `*_EAGER` rungs of a same-binary A/B. Rerouting +any of them changes what the reference computes, which deletes the comparison the +gate performs; AGENTS.md names that failure directly — never make a red gate green +by widening an assertion's scope. A documented opt-in helper that picks by shape +was also rejected as insufficient on its own: it helps an author who already knows +the fast rungs exist, and #1544's defect is precisely the author who does not. + +The checker inverts that. It cannot change any caller's numerics, because it runs +no code; and it fires on the one population that matters, an author naming the +naive kernel without saying why. + +**D2 — narrow the bound, do not opt in to a larger cap.** `cudaFuncSetAttribute` +with `cudaFuncAttributeMaxDynamicSharedMemorySize` would make the advertised 256 +true on some devices and NOT on others: head_dim 256 in f32 needs 128 KiB, above +the opt-in per-block cap of the consumer Blackwell parts this project gates on, so +the opt-in call itself can fail and the contract would still be a lie at the +widest advertised width. Narrowing is device-independent arithmetic, provable +here, and is a strict improvement for every caller: a head_dim that launches today +still launches, and one that does not now fails with a message naming the rung +that works instead of an opaque CUDA launch error from a later +`cudaGetLastError`. + +That paragraph originally added "and it cannot be verified without a device, and +this row has no lease", which made the rejection read as a convenience. It is not +one, and the number is now on the record: **GB10's queried opt-in ceiling is +101,376 bytes**, measured while #1557 was reviewed. head_dim 256 in f32 wants +131,072. Opting in therefore CANNOT make the advertised 256 true for f32 on the +part this project gates on — the raise buys nothing at the width that motivated +it, and a caller at that width would have gone on falling back silently without +ever launching. Narrowing beats raising here on a measurement rather than on a +preference, which is why this row lands first and unchanged rather than +reconciling onto a cap-raising change. + +Opting in remains available later as a widening for bf16 widths above 192, owned +by nobody today because no live caller needs head_dim above the honest bound +through this op. + +**D3 — refuse rather than silently fall back to `AttentionDenseFast`.** +`AttentionDenseFa2KernelCuda` DOES fall through to `AttentionDenseFlash` +(`cuda_ops.cu:3395-3408`), so a silent step-down has precedent here. It was +rejected anyway. `AttentionDenseFast` re-reads K and V from global memory once per +(query, head) — that is the exact redundancy `AttentionDenseFlash` exists to +remove — so the fallback is a real, unannounced slowdown, which is #1544's disease +in miniature. This row's whole subject is that a rung change must be a declared +choice. A refusal naming `AttentionDenseFast` gives the caller the same +information and lets them make it. + +**D4 — the allowlist is not a lock.** AGENTS.md forbids a record surface every +pull request must write. `scripts/attention-rung-allowlist.txt` is written only by +a change that ADDS an unmarked naive-attention site, which is the event the +checker exists to make deliberate, and it is emptied by the rows already in +flight. Two concurrent removals of different lines merge. The primary record is +still per-site and in-file: the marker comment lives beside the call it explains, +so the ordinary case touches no shared file at all. + +**D5 — the marker's reason is checked for presence, not for truth.** The checker +requires a non-trivial reason string and cannot judge it. That is the same floor +`check-fusion-consistency.py` sets with its allowlist reasons. A reviewer judges +the reason; the gate only guarantees one was written. + +**D6 — the detected population is one literal spelling, and that is stated.** +The scan matches `vt::Attention(` in a model `.cpp` or `.h`. Four spellings reach +the same kernel and are NOT detected — a `using vt::Attention;` plus a bare call, 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 none is how attention is called here, so this is a stated bound and +not a live hole. Widening the regex was rejected: dropping the `vt::` prefix makes +every fast rung a site, which is D1's failure mode again, and no regex reaches a +function pointer at all. Closing it needs a compiler-side population (the op +registry, or clang tooling over the real translation unit), which is a different +instrument and not this row's scope. The checker's docstring says so, so a green +reads as "no unmarked `vt::Attention(` call" and never as "no model is naive". + +**D7 — the allowlist's stem set is pinned by a test, in another file.** D4 says the +CHECKER never forces a removing row to edit the allowlist, and that is still true. +`test_allowlist_holds_only_the_in_flight_stems` does force it: the expected set is +pinned, so adding or deleting a stem reds that case until the test is updated in +the same change. That is the intended shape for a parking lot — growth must be a +review decision — but it is a coupling a reader of the allowlist alone would not +see, so the allowlist header and the checker docstring both name the test. + +**R1 — the launcher refusal is not executed on this box.** The pure arithmetic is +tested and mutated here, but nothing on a CPU-only box proves the launcher CALLS +it. A reviewer's reachability mutation for that leg needs a CUDA device. Stated, +not hidden; see `## Owed`. + +**R2 — the 48 KiB constant.** `49152` is CUDA's default dynamic shared-memory cap +on every architecture this project supports, and the bound is INCLUSIVE. That +matters in one direction only: head_dim 192 in bf16 sits exactly at 49152 and +launches today, so an exclusive bound would refuse work that currently runs, which +would be a regression rather than a repair. + +## Owed + +- [#1573](https://github.com/mudler/vllm.cpp/issues/1573) — run the CUDA head-dim + refusal case for `AttentionDenseFlash` on a leased device, and mutate the + launcher's bound call to prove the case reaches it. PENDING a GPU lease. This + stays owed after the merge: the CPU cases pin the arithmetic and never that the + launcher calls it, and no lease was available for the whole branch. D2's + 101,376-byte GB10 ceiling is a queried device value and does not discharge it, + because it bounds what an opt-in could buy rather than proving the refusal + executes. + +- [#1629](https://github.com/mudler/vllm.cpp/issues/1629) — DISCHARGED IN THIS + ROW, and listed here because AGENTS.md wants the index row, this spec and the + pull request body to agree, and this section is where this spec links its + issues. Two drift locks in `tests/scripts/test_check_attention_rung_consistency.py` + stored counts of files they do not own, so the three rows the attention-rung + allowlist exists to unblock had no green path: `test_the_population_is_not_empty` + asserted `>= 9` against a tree holding exactly 9 `vt::Attention(` sites, and + `assertGreater(excused, 0)` required the shipped allowlist to stay non-empty + forever. Both are repaired here, and neither by lowering a number, which is the + mute-switch failure: the floor is now `>= 1`, which detects only a broken + scanner and leaves the rename guard to + `test_the_six_deliberate_sites_carry_a_marker`, which pins six sites by name; a + new case proves the stem a red message names is a real source file; and the + excused counter is pinned by two synthetic cases that build their own + allowlisted population. Nothing is owed after the merge, so this entry is the + link and not a debt. + +- [#1631](https://github.com/mudler/vllm.cpp/issues/1631) — teach + `scripts/check-pr-size.py` to tell a comment-only or docstring-only diff to a + `governance_checker` from a semantic one, so a measurably false comment in a + checker can be corrected on its own. `check-pr-size.py:170` classifies every + `scripts/check-*.py` and `.sh` as a governance checker and `change_errors` then + requires paired test evidence that goes red against the BASE checker, which a + semantically identical diff cannot produce by construction. OWED after the + merge, and it is why three comments in + `scripts/check-attention-rung-consistency.py` ship unrepaired beside a suite + that was repaired for #1629: `:58-61` gives the wrong reason for not widening + the regex, `:93-96` attributes the exclusion of the fast rungs to the `\b` + rather than to the trailing `\(`, and `:252-255` denies an equality that holds + on this tree (9 sites, 6 marked, 3 excused). It cannot be fixed here, because + changing what the gate accepts is what AGENTS.md `## Changing the rules or a + checker` routes to its own row, spec and red-before evidence; attaching the + correction to an unrelated semantic change is the alternative that section + exists to refuse. A candidate patch is parked on the issue. + +## Now + +The change is written, CPU-gated and through one fresh scoped review, whose +findings are repaired here: the new checker registers its disabled creation- +mutation stub in `check-pr-size.py` (measured 31 of 31 cases red under the stub); +the kernel's register blocking is hoisted to `kFlashMaxPerLane` at file scope so +the `static_assert` reads the constant the kernel uses instead of the literal `8`; +the launcher's comment no longer claims the guard and the shared-memory request +come from one function; the checker states the four spellings it does not detect +and reports how many sites are unmarked and excused; and the kernel-matrix cell no +longer stores this suite's case count. The single owed leg above still needs +whoever next holds a lease. diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index c1f0cad49..0afaa1e3a 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -243,6 +243,10 @@ jobs: run: | python3 scripts/check-fusion-consistency.py python3 tests/scripts/test_check_fusion_consistency.py + - name: A model on the naive attention kernel says why (#1544) + run: | + python3 scripts/check-attention-rung-consistency.py + python3 tests/scripts/test_check_attention_rung_consistency.py - name: Structural checkers ignore text the compiler never sees run: | python3 tests/scripts/test_checker_text.py diff --git a/include/vt/ops.h b/include/vt/ops.h index aa737b63f..89bdd7be3 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -3320,6 +3320,46 @@ void AttentionDenseFast(Queue& q, Tensor& out, const Tensor& query, const Tensor void AttentionDenseFlash(Queue& q, Tensor& out, const Tensor& query, const Tensor& key, const Tensor& value, const AttentionArgs& args); +// The head_dim domain AttentionDenseFlash can actually LAUNCH, as arithmetic a +// host without a GPU can execute and a test can pin (#1544). +// +// The tiling is what bounds it: each CTA stages `kAttentionDenseFlashTileCols` +// columns of BOTH K and V in dynamic shared memory, so it asks the driver for +// `2 * cols * head_dim * sizeof(input element)` bytes. There is no +// `cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, ...)` +// anywhere in src/vt/cuda/, so that request is capped at CUDA's DEFAULT 48 KiB per +// block on every architecture — not at the 256 the kernel's register blocking +// allows, which is what this op used to advertise on its own. head_dim 256 in bf16 +// wants 64 KiB and in f32 wants 128 KiB; both are refused by the driver, and the +// caller used to learn that from a bare launch error one `cudaGetLastError` later. +// +// This mirrors what vLLM makes every backend declare — `get_supported_head_sizes` +// / `supports_head_size`, vllm/v1/attention/backend.py:155-163 @ 555967922 — where +// the domain is consulted BEFORE dispatch rather than discovered by launching. +// AttentionDenseFast (kAttentionDenseFast) uses NO shared memory and is the rung +// that serves head_dim above the bound below. +inline constexpr int64_t kAttentionDenseFlashTileCols = 64; +// The register blocking: 8 elements per lane across 32 lanes. +inline constexpr int64_t kAttentionDenseMaxHeadDim = 256; +// CUDA's default per-block dynamic shared-memory cap, and it is INCLUSIVE. That +// direction matters: head_dim 192 in bf16 lands on exactly 49152 bytes and launches +// today, so an exclusive bound would refuse work that currently runs. +inline constexpr int64_t kCudaDefaultDynamicSmemBytes = 49152; + +// Bytes of dynamic shared memory one CTA requests for `head_dim` at `elem_size`. +constexpr int64_t AttentionDenseFlashSmemBytes(int64_t head_dim, int64_t elem_size) { + return 2 * kAttentionDenseFlashTileCols * head_dim * elem_size; +} + +// The largest head_dim AttentionDenseFlash can launch for an input element size. +// bf16 -> 192, f32 -> 96. +constexpr int64_t AttentionDenseFlashMaxHeadDim(int64_t elem_size) { + if (elem_size <= 0) return 0; // no element size, no admissible head_dim + const int64_t by_smem = + kCudaDefaultDynamicSmemBytes / (2 * kAttentionDenseFlashTileCols * elem_size); + return by_smem < kAttentionDenseMaxHeadDim ? by_smem : kAttentionDenseMaxHeadDim; +} + // Same contract as AttentionDenseFlash, but the CUDA impl runs the VENDORED // FlashAttention-2 forward (src/vt/cuda/flash_attn/) on its tensor cores instead of a // scalar per-warp recurrence — the kernel vLLM itself dispatches for dense non-causal diff --git a/scripts/agent-preflight.sh b/scripts/agent-preflight.sh index a094cf60e..7b18b25cb 100755 --- a/scripts/agent-preflight.sh +++ b/scripts/agent-preflight.sh @@ -99,6 +99,7 @@ CHECKERS=( check-supported-models check-env-doc check-fusion-consistency + check-attention-rung-consistency check-fp4-resident-consistency check-cuda-op-arch-gate check-runner-routing-consistency @@ -149,6 +150,7 @@ SUITES=( test_check_env_doc test_checker_text test_check_fusion_consistency + test_check_attention_rung_consistency test_check_fp4_resident_consistency test_check_cuda_op_arch_gate test_check_runner_routing_consistency diff --git a/scripts/attention-rung-allowlist.txt b/scripts/attention-rung-allowlist.txt new file mode 100644 index 000000000..22041a1e5 --- /dev/null +++ b/scripts/attention-rung-allowlist.txt @@ -0,0 +1,36 @@ +# Attention-rung allowlist — model TUs whose `vt::Attention` call (the naive, +# correctness-grade kernel) is being REMOVED by a row already in flight, so +# scripts/check-attention-rung-consistency.py stays green without this change +# editing the very lines those rows replace. +# +# This is NOT the place to park a naive call you intend to keep. A deliberate +# call records its reason IN THE FILE, beside the call: +# +# // VT-ATTN-NAIVE: +# +# which is why the six deliberate sites (whisper_audio, qwen3_vl_vision, +# kimi_linear_device, qwen3_5, nemotron_h, nemotron_h_device) are NOT listed here. +# Full reasoning: .agents/specs/attention-rung-visibility.md D4. +# +# Each entry is a model-file stem (src/vllm/model_executor/models/.cpp) and +# names the issue that owns its removal. Deleting a stem here, after that row +# routes the call to a fast rung, is the enforcement closing. The checker reports +# a stem whose sites are gone or now marked as STALE and does NOT fail on it, so +# the removing row is free to leave the deletion to whoever runs preflight next. +# +# The CHECKER does not fail on it. A TEST does, and it lives in another file: +# tests/scripts/test_check_attention_rung_consistency.py +# ::ShippedTreeTests::test_allowlist_holds_only_the_in_flight_stems pins this set +# exactly, so adding OR deleting a stem here reds that case until the expected set +# is updated in the same change. That is deliberate -- growth of a parking lot must +# be a review decision and not a silent edit -- but it means "the removing row owes +# this file nothing" is true of the checker only. A row that deletes its stem edits +# that test too. + +# --- IN FLIGHT: the naive call is the defect, and another row is removing it --- +muse_glimmer_vision # 50 layers, H=16, head_dim=96, non-causal, sole path, no knob. + # Issue #1545; the fix is vt::AttentionDenseFlash (24 KB of + # shared memory at head_dim 96, inside the honest bound). +ltx2 # DiT self-attention; 47.84 s measured per forward (#1544). + # LTX-2.5 routing row in flight. +ltx2_device # the device arm of the same DiT forward, same row. diff --git a/scripts/check-attention-rung-consistency.py b/scripts/check-attention-rung-consistency.py new file mode 100755 index 000000000..b29629571 --- /dev/null +++ b/scripts/check-attention-rung-consistency.py @@ -0,0 +1,273 @@ +#!/usr/bin/env python3 +r"""Fail if a model names the NAIVE attention kernel without saying why. + +`vt::Attention` resolves `OpId::kAttention` (`src/vt/ops.cpp`) straight to +`AttentionKernel` (`src/vt/cuda/cuda_ops.cu`), self-described there as +"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, K and V re-read from +global once per (query, head). Which kernel a model gets is decided by the C++ +function name its author typed, and nothing ever routes `kAttention` up. A model +whose author never heard of `vt::AttentionDenseFlash` / `vt::AttentionDenseFast` / +`vt::AttentionDenseFa2` therefore pays up to ~500x with correct output, no warning +and no gate (issue #1544; measured at 47.84 s for one LTX-2.5 DiT forward). + +The freeze is NOT the defect. `kAttention` is frozen so text decode stays +byte-identical, and six model sites use the naive kernel deliberately: as the +reference arm of a numeric gate, or as the `VT_*_EAGER` rung of a same-binary A/B. +The defect is that nothing distinguishes those six from an author who simply did +not know. So this checker does not reroute anything and cannot: it requires the +CHOICE to be recorded next to the call. + +The record is a marker comment on the call line or in the 20 lines above it: + + // VT-ATTN-NAIVE: reference arm of the paged/dense equivalence gate; the fast + // rungs are not bit-identical to this one, so rerouting deletes the golden. + vt::Attention(q, out, qq, kk, vv, args); + +`whisper_audio.cpp` and `qwen3_vl_vision.cpp` are the intended pattern (#1544): +both DEFAULT to a fast rung and expose the naive one behind an env knob. + +Why a marker and not one shared registry: AGENTS.md forbids a record surface every +pull request must write. The marker lives in the file that owns the call, so the +ordinary case touches no shared file. `scripts/attention-rung-allowlist.txt` +carries only stems whose naive call is being REMOVED by a row already in flight — +editing the very lines those changes replace would conflict for no gain. An +allowlisted stem whose sites are all marked or gone is reported as STALE and does +NOT fail HERE, so the CHECKER never forces the row that cleans it up to edit this +file. A test does: the expected stem set is pinned in +tests/scripts/test_check_attention_rung_consistency.py, which the allowlist's own +header says, so growth of that file stays a review decision. + +Text the compiler never sees is not a call site: the scan runs over +`scripts/checker_text.py::normalize_source`, so a commented-out, `#if 0`-ed or +`if (false)`-ed `vt::Attention(` is a deletion here, exactly as it is to nvcc. +That normalization is position-preserving, so every reported `file:line` still +describes the original file. + +What this checker DETECTS, stated as a limit rather than implied by a green: +one literal spelling, `vt::Attention(`, in a model `.cpp` or `.h`. Four spellings +reach the same kernel and are NOT detected, each verified to leave the checker +green with a live unmarked call: + + using vt::Attention; then a bare Attention(...) + namespace vv = vt; then vv::Attention(...) + #define MUT_ATTN vt::Attention + auto* fn = &vt::Attention; then a call through `fn` + +None exists in this tree, and the repository does not write attention calls that +way, so this is a stated bound and not a live hole. Widening the regex is not the +repair: `\bAttention\s*\(` also matches every fast rung's suffix-free form and +would demand a marker beside exactly the calls this checker wants people to make, +and no regex reaches a call through a function pointer at all. What closes it is a +compiler-side population — the CUDA op registry, or a clang tooling pass over the +real translation unit — which is a different instrument, not a longer pattern. A +green here therefore means "no unmarked `vt::Attention(` call", never "no model is +on the naive rung". + +The validation logic is pure functions (`scan_file`, `drift_sites`, +`stale_allowlist_entries`) so it is unit- and mutation-testable +(tests/scripts/test_check_attention_rung_consistency.py), mirroring +check-fusion-consistency.py. +""" + +from __future__ import annotations + +import re +import sys +from pathlib import Path + +ROOT = Path(__file__).resolve().parents[1] +sys.path.insert(0, str(ROOT / "scripts")) + +from checker_text import normalize_source # noqa: E402 + +# Both halves of a model's sources. A header is scanned too, because a call moved +# into an inline function or a template there would otherwise leave the checker +# green — the population is what makes a green meaningful. +MODEL_DIRS = ( + ROOT / "src/vllm/model_executor/models", + ROOT / "include/vllm/model_executor/models", +) +ALLOWLIST = ROOT / "scripts/attention-rung-allowlist.txt" + +# `vt::Attention(` and nothing else. The word boundary is load-bearing: without it +# this matches `vt::AttentionDenseFlash(`, `vt::AttentionDenseFast(`, +# `vt::AttentionDenseFa2(` and `vt::AttentionCross(` — every FAST rung — and the +# checker would demand a marker beside exactly the calls it wants people to make. +_NAIVE_CALL = re.compile(r"\bvt::Attention\s*\(") + +# The marker, in a `//` comment. A reason is required after the colon. +_MARKER = re.compile(r"//.*\bVT-ATTN-NAIVE:\s*(\S.*?)\s*$") + +# How far above a call the marker may sit. A call reached through several lines of +# tensor-view setup still reads as one statement to a human, and forcing the +# marker onto the call line itself would push it past any sane column. +MARKER_WINDOW_LINES = 20 + +# A reason must be long enough to BE one. This is a floor against `// VT-ATTN-NAIVE: x` +# and nothing more; the checker cannot judge whether a reason is true, and does not +# try. That is a reviewer's job, exactly as it is for the reasons on +# scripts/fusion-consistency-allowlist.txt. +MIN_REASON_CHARS = 16 + + +def marker_reason(line: str) -> str | None: + """The reason recorded by a marker comment on this line, or None.""" + match = _MARKER.search(line) + return match.group(1) if match is not None else None + + +def has_marker(lines: list[str], call_line: int) -> bool: + """True if a marker with a substantive reason covers the call at `call_line`. + + `call_line` is 1-based. The window is the call line itself and the + MARKER_WINDOW_LINES lines above it. + """ + first = max(1, call_line - MARKER_WINDOW_LINES) + for number in range(first, call_line + 1): + reason = marker_reason(lines[number - 1]) + if reason is not None and len(reason) >= MIN_REASON_CHARS: + return True + return False + + +def scan_file(text: str) -> list[tuple[int, bool]]: + """Every live `vt::Attention(` site in one translation unit. + + Returns (1-based line number, marker_present) per site, in source order. + """ + live = normalize_source(text) + raw_lines = text.splitlines() + out: list[tuple[int, bool]] = [] + for match in _NAIVE_CALL.finditer(live): + line = live.count("\n", 0, match.start()) + 1 + out.append((line, has_marker(raw_lines, line))) + return out + + +def scan_models(model_dirs=MODEL_DIRS) -> dict[str, list[tuple[int, bool]]]: + """Map repo-relative model source path -> its `vt::Attention(` sites. + + Keyed on the PATH, not the stem: `ltx2.cpp` and `ltx2.h` share a stem and would + otherwise overwrite each other, reporting one file's sites under the other's + name. The allowlist still matches on the stem, so one entry covers a model's + whole translation unit. + """ + out: dict[str, list[tuple[int, bool]]] = {} + for models_dir in model_dirs: + if not models_dir.is_dir(): + continue + for pattern in ("*.cpp", "*.h"): + for path in sorted(models_dir.glob(pattern)): + sites = scan_file(path.read_text(encoding="utf-8", errors="ignore")) + if sites: + out[str(path.relative_to(ROOT))] = sites + return out + + +def allowlisted_names(text: str) -> set[str]: + """Model stems accepted as in-flight (one per line, # comments ignored) — + mirrors check-fusion-consistency.py.""" + names: set[str] = set() + for line in text.splitlines(): + line = line.split("#", 1)[0].strip() + if line: + names.add(line) + return names + + +def drift_sites( + scanned: dict[str, list[tuple[int, bool]]], allowlisted: set[str] +) -> list[tuple[str, int]]: + """(path, line) for every unmarked naive-attention call in a file whose stem is + not allowlisted. Empty == the check passes.""" + out: list[tuple[str, int]] = [] + for path in sorted(scanned): + if Path(path).stem in allowlisted: + continue + out.extend((path, line) for line, marked in scanned[path] if not marked) + return out + + +def stale_allowlist_entries( + scanned: dict[str, list[tuple[int, bool]]], allowlisted: set[str] +) -> list[str]: + """Allowlisted stems that would pass on their own merit — their naive calls are + gone or now carry a marker. Reported, never fatal: the row that removes the + call must not be forced to edit this file to stay green.""" + return sorted( + stem + for stem in allowlisted + if all( + marked + for path, sites in scanned.items() + if Path(path).stem == stem + for _, marked in sites + ) + ) + + +def main() -> int: + scanned = scan_models() + allowlisted = ( + allowlisted_names(ALLOWLIST.read_text(encoding="utf-8")) + if ALLOWLIST.exists() + else set() + ) + drift = drift_sites(scanned, allowlisted) + + for stem in stale_allowlist_entries(scanned, allowlisted): + print( + f"STALE (not a failure): {stem} is on " + "scripts/attention-rung-allowlist.txt but no longer has an unmarked " + "vt::Attention call. Delete its entry." + ) + + if drift: + print( + "ERROR: model forward(s) call vt::Attention — the naive, " + "correctness-grade attention kernel (up to ~500x the cost of " + "vt::AttentionDenseFlash; issue #1544) — with no recorded reason:", + file=sys.stderr, + ) + for path, line in drift: + print(f" - {path}:{line}", file=sys.stderr) + print( + "If the fast rung is what you wanted, call vt::AttentionDenseFlash " + "(shared-memory tiled), vt::AttentionDenseFast (warp-per-query, no " + "shared memory) or vt::AttentionDenseFa2 (bf16 head_dim 64 " + "non-causal tensor cores). If the NAIVE kernel is what you meant — a " + "reference arm, or the eager rung of a same-binary A/B — record that " + "on the call line or within " + f"{MARKER_WINDOW_LINES} lines above it:\n" + " // VT-ATTN-NAIVE: \n" + "See src/vllm/model_executor/models/whisper_audio.cpp and " + "qwen3_vl_vision.cpp for the intended pattern.", + file=sys.stderr, + ) + return 1 + + sites = sum(len(v) for v in scanned.values()) + marked = sum(1 for v in scanned.values() for _, m in v if m) + # The number a reader of a green actually needs: sites that carry NO reason and + # pass only because their stem is allowlisted. `sites - marked` is not it, because + # a marked call inside an allowlisted file counts in `marked`. This is the debt the + # green is hiding, so it is printed even when it is zero. + excused = sum( + 1 + for path, sites_in_file in scanned.items() + if Path(path).stem in allowlisted + for _, m in sites_in_file + if not m + ) + print( + f"OK (attention rung): {sites} vt::Attention call site(s) in " + f"{len(scanned)} model source file(s); {marked} carry a recorded reason, " + f"{excused} unmarked and excused by " + f"{len(allowlisted)} allowlisted in-flight stem(s)." + ) + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/scripts/check-pr-size.py b/scripts/check-pr-size.py index 33ec85bff..b153cc713 100755 --- a/scripts/check-pr-size.py +++ b/scripts/check-pr-size.py @@ -365,6 +365,13 @@ # code of 1 or an examined count out of the report goes red, which is what # makes the stub a mutation rather than a weaker checker. "scripts/check-conflict-markers.py": DISABLED_CREATION_CHECKER, + # KERNEL-ATTN-DENSE-FLASH (#1544). Created in this range, so there is no BASE + # version to mutate. Its suite loads the checker as a module at import time and + # then calls scan_file / has_marker / drift_sites / stale_allowlist_entries / + # main on it, none of which the stub defines, so every case that touches the + # checker raises AttributeError. Measured: 31 of 31 red under the stub, because + # the suite has no case that passes without calling into the checker at all. + "scripts/check-attention-rung-consistency.py": DISABLED_CREATION_CHECKER, } SELF_CHECKER = "scripts/check-pr-size.py" EVIDENCE_TIMEOUT_SECONDS = 120 diff --git a/src/vllm/model_executor/models/kimi_linear_device.cpp b/src/vllm/model_executor/models/kimi_linear_device.cpp index 8faf25b70..84ac4358a 100644 --- a/src/vllm/model_executor/models/kimi_linear_device.cpp +++ b/src/vllm/model_executor/models/kimi_linear_device.cpp @@ -595,6 +595,13 @@ DBuf MlaAttnCoreDevice(const Dev& d, DBuf& dq, DBuf& dkv, DBuf& dkpe, Tensor query = MakeTensor(dq.ptr(), DType::kF32, d.q.device, {T, nah, qk}); DBuf attn(d, DType::kF32, {T, nah, qk}); const float scale = static_cast(std::pow(static_cast(qk), -0.5)); + // VT-ATTN-NAIVE: the whole point of this arm is that the attention core runs on + // the SAME f32 online max-subtracted softmax as the host reference, as the + // MlaAttnCoreDevice design note at the top of this file records, and it is + // behind VT_KIMI_DEVICE_MLA, default OFF, recorded there as a measured + // negative (4.24 -> 3.89 tok/s). No fast rung is available here either: the + // padded head_dim is qk=192 in f32, and vt::AttentionDenseFlash would need 96 KB + // of dynamic shared memory, twice CUDA's default cap (#1544, cuda_ops.cu). vt::Attention(d.q, attn.t(), query, key.t(), val.t(), vt::AttentionArgs{scale, true}); // slice out[:, :, :vh] -> [T, nah*vh] (the pad-V tail is 0 by construction). diff --git a/src/vllm/model_executor/models/nemotron_h.cpp b/src/vllm/model_executor/models/nemotron_h.cpp index aac752bac..df3edd935 100644 --- a/src/vllm/model_executor/models/nemotron_h.cpp +++ b/src/vllm/model_executor/models/nemotron_h.cpp @@ -668,6 +668,11 @@ std::vector NemotronHAttentionMixer(const NemotronHAttentionWeights& w, // `self.scaling = self.head_dim**-0.5` (nemotron_h.py:440). args.scale = static_cast(1.0 / std::sqrt(static_cast(Dh))); args.causal = true; + // VT-ATTN-NAIVE: the HOST reference arm. `NemotronHAttnBlock` in + // nemotron_h_device.cpp is the device arm, and the equivalence gate between + // them holds only while both run the same kernel — the fast rungs are NOT + // bit-identical to this one. Changing one side reroutes the comparison + // rather than the model (#1544). vt::Attention(queue, ot, qt, kt, vt_, args); } diff --git a/src/vllm/model_executor/models/nemotron_h_device.cpp b/src/vllm/model_executor/models/nemotron_h_device.cpp index 4455998d9..5e83a0a57 100644 --- a/src/vllm/model_executor/models/nemotron_h_device.cpp +++ b/src/vllm/model_executor/models/nemotron_h_device.cpp @@ -327,6 +327,11 @@ DBuf NemotronHAttnBlock(Dev d, const NemotronHAttentionWeights& w, // narrowing, so the two arms feed `vt::Attention` bit-identical scales. args.scale = static_cast(1.0 / std::sqrt(static_cast(Dh))); args.causal = true; + // VT-ATTN-NAIVE: the DEVICE half of the pair above. `NemotronHAttentionMixer` + // in nemotron_h.cpp is the host reference, and this arm exists to be + // bit-comparable to it — same scale, same rounding, same kernel. A fast rung + // here would make the two arms measure different things instead of the same + // one (#1544). vt::Attention(d.q, attn.t(), qt, kt, vt_, args); } diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index a5996ec9b..9f32a034d 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -5276,6 +5276,11 @@ DBuf FullAttnBlock(Dev d, const FullAttnLayerWeights& w, const HfConfig& cfg, } DBuf dattn(d, DType::kF32, {T, Hq, Dh}); const float scale = 1.0F / std::sqrt(SizeF(Dh)); + // VT-ATTN-NAIVE: the REFERENCE (non-paged) dense arm, as the comment on the V + // upcast above already says. Production decode runs `FullAttnBlockPaged`, which + // replaces this call with vt::ReshapeAndCache + vt::PagedAttention; this arm is + // what that path is compared against, so a rung change here moves the golden + // rather than the shipping kernel (#1544). vt::Attention(d.q, dattn.t(), qn3, kn3, v3, vt::AttentionArgs{scale, true}); // Sigmoid output gate on the raw gate split, folded into the o_proj activation diff --git a/src/vllm/model_executor/models/qwen3_vl_vision.cpp b/src/vllm/model_executor/models/qwen3_vl_vision.cpp index 7ae12dae8..7992fbb0d 100644 --- a/src/vllm/model_executor/models/qwen3_vl_vision.cpp +++ b/src/vllm/model_executor/models/qwen3_vl_vision.cpp @@ -523,6 +523,10 @@ std::vector Qwen3VLVisionForward(const std::vector& pixel_value // grid_t==1 == single 784-patch window). Default flash-tiled (byte-identical // to warp; see vis_attn above). kAttention (text/audio) is untouched. const vt::AttentionArgs aargs{scale, /*causal=*/false}; + // VT-ATTN-NAIVE: the EAGER rung of the same-binary A/B above, reachable + // only with VT_QWEN3VL_ATTN_EAGER=1, which nothing in the tree sets. The + // default is the flash-tiled rung two branches down. Rerouting this arm + // would delete the baseline the other two are measured against (#1544). if (vis_attn == 0) vt::Attention(q, aof, qf, kf, vf, aargs); else if (vis_attn == 1) diff --git a/src/vllm/model_executor/models/whisper_audio.cpp b/src/vllm/model_executor/models/whisper_audio.cpp index f60184596..5d92b4c18 100644 --- a/src/vllm/model_executor/models/whisper_audio.cpp +++ b/src/vllm/model_executor/models/whisper_audio.cpp @@ -320,6 +320,10 @@ std::vector WhisperAudioEncoderForward(const std::vector& input_fe if (f2 != nullptr && f2[0] == '1') return 3; // FA-2 tensor cores (opt-in) return 2; // flash-tiled (default, byte-exact) }(); + // VT-ATTN-NAIVE: the EAGER rung of the same-binary A/B above, reachable only + // with VT_WHISPER_ENC_EAGER=1, which nothing in the tree sets. The default is + // the flash-tiled rung two branches down. Rerouting this arm would delete the + // baseline the other three are measured against (#1544). if (enc_attn == 0) vt::Attention(q, ao.tensor(), q3, k3, v3, vt::AttentionArgs{scale, /*causal=*/false}); else if (enc_attn == 1) diff --git a/src/vt/cuda/cuda_ops.cu b/src/vt/cuda/cuda_ops.cu index f6337e4c6..0228e53fe 100644 --- a/src/vt/cuda/cuda_ops.cu +++ b/src/vt/cuda/cuda_ops.cu @@ -3234,12 +3234,24 @@ void AttentionDenseFastKernelCuda(Queue& q, Tensor& out, const Tensor& query, co // untouched. One q-head per CTA (all warps share the same GQA kv-head g). constexpr int kFlashBr = 16; // query-warps per CTA (= K/V global-read reuse factor) constexpr int kFlashBc = 64; // key/value columns streamed per shared-memory tile +// The register blocking: head_dim elements each of the 32 lanes holds. It lives at +// file scope rather than inside the kernel body precisely so the static_assert below +// can READ it -- a per-kernel local is invisible here, and asserting the literal 8 +// instead would compare the header's constant against a number nothing else uses. +constexpr int kFlashMaxPerLane = 8; // head_dim up to 8 * 32 +// The head_dim bound this kernel advertises is computed in include/vt/ops.h so a box +// with no GPU can execute it. These tie the two together: change the tile width or the +// register blocking here and the arithmetic there stops describing this kernel, which +// is how the op came to advertise a head_dim it could not launch (#1544). +static_assert(kFlashBc == kAttentionDenseFlashTileCols, + "AttentionDenseFlashSmemBytes must use this kernel's tile width"); +static_assert(kFlashMaxPerLane * 32 == kAttentionDenseMaxHeadDim, + "kFlashMaxPerLane * warp size must equal the advertised register bound"); template __global__ void AttentionDenseFlashKernel(Tout* out, const Tin* query, const Tin* key, const Tin* value, int64_t hq, int64_t hk, int64_t d, int64_t t, float scale, bool causal) { - constexpr int kMaxPerLane = 8; // head_dim up to 256 extern __shared__ __align__(16) char flash_smem[]; Tin* sK = reinterpret_cast(flash_smem); Tin* sV = sK + static_cast(kFlashBc) * d; @@ -3254,10 +3266,10 @@ __global__ void AttentionDenseFlashKernel(Tout* out, const Tin* query, const Tin const int npl = static_cast((d + 31) / 32); // head_dim elements this lane owns // This warp's query row, in registers (identical layout to AttentionWarpKernel). - float qreg[kMaxPerLane]; - float acc[kMaxPerLane]; + float qreg[kFlashMaxPerLane]; + float acc[kFlashMaxPerLane]; #pragma unroll - for (int k = 0; k < kMaxPerLane; ++k) { + for (int k = 0; k < kFlashMaxPerLane; ++k) { qreg[k] = 0.0f; acc[k] = 0.0f; } @@ -3298,7 +3310,7 @@ __global__ void AttentionDenseFlashKernel(Tout* out, const Tin* query, const Tin const int64_t base = j * d; float part = 0.0f; #pragma unroll - for (int k = 0; k < kMaxPerLane; ++k) { + for (int k = 0; k < kFlashMaxPerLane; ++k) { const int e = lane + 32 * k; if (k < npl && e < d) part += qreg[k] * Load(sK, base + e); } @@ -3309,7 +3321,7 @@ __global__ void AttentionDenseFlashKernel(Tout* out, const Tin* query, const Tin const float corr = expf(m - m_new); const float p = expf(s - m_new); #pragma unroll - for (int k = 0; k < kMaxPerLane; ++k) { + for (int k = 0; k < kFlashMaxPerLane; ++k) { const int e = lane + 32 * k; if (k < npl && e < d) acc[k] = acc[k] * corr + p * Load(sV, base + e); } @@ -3332,10 +3344,40 @@ void LaunchAttentionDenseFlash(cudaStream_t s, Tensor& out, const Tensor& query, const int64_t t = query.shape[0], hq = query.shape[1], d = query.shape[2]; const int64_t hk = key.shape[1]; if (t == 0 || hq == 0 || d == 0) return; - VT_CHECK(d <= 256, "cuda attention-dense-flash: head_dim <= 256 only"); + // The honest bound, not the register bound. `kFlashMaxPerLane` allows head_dim 256, but + // the K/V tile below asks for `2*kFlashBc*d*sizeof(Tin)` bytes of DYNAMIC shared + // memory, and no `cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemory + // Size, ...)` exists anywhere in src/vt/cuda/ — so the driver caps the request at + // its default 48 KiB and the real ceiling is 192 (bf16) / 96 (f32). This op used to + // advertise `d <= 256` and hand a wider caller a bare launch failure from the + // `cudaGetLastError` at the bottom of this function, naming nothing it could do + // about it (#1544). REFUSING here rather than falling back to AttentionDenseFast is + // deliberate: that rung re-reads K and V from global once per (query, head), which + // is the exact redundancy this kernel exists to remove, so taking it silently would + // be an unannounced slowdown. Name it and let the caller choose. + const int64_t dmax = AttentionDenseFlashMaxHeadDim(static_cast(sizeof(Tin))); + VT_CHECK(d <= dmax, + std::string("cuda attention-dense-flash: head_dim ") + std::to_string(d) + + " needs " + + std::to_string(AttentionDenseFlashSmemBytes( + d, static_cast(sizeof(Tin)))) + + " bytes of dynamic shared memory, over CUDA's default cap of " + + std::to_string(kCudaDefaultDynamicSmemBytes) + + "; this kernel serves head_dim <= " + std::to_string(dmax) + + " for this input dtype. Use vt::AttentionDenseFast, which uses no " + "shared memory and serves head_dim <= " + + std::to_string(kAttentionDenseMaxHeadDim)); const unsigned nblk = static_cast((t + kFlashBr - 1) / kFlashBr); const dim3 grid(nblk, static_cast(hq)); - const size_t shmem = static_cast(2) * kFlashBc * d * sizeof(Tin); // sK + sV + // The request the guard above admitted. They are two functions, and + // AttentionDenseFlashMaxHeadDim RE-DERIVES the division rather than inverting + // AttentionDenseFlashSmemBytes, so agreement is a property to be tested, not one + // the code makes structural. tests/vt/test_ops_attention.cpp pins it in both + // directions at the inclusive edge; mutating the `2 *` in SmemBytes to `3 *` turns + // those cases RED, which is what keeps the two halves honest. Open-coding the byte + // count here instead is how this contract drifted the first time. + const size_t shmem = static_cast( + AttentionDenseFlashSmemBytes(d, static_cast(sizeof(Tin)))); // sK + sV switch (out.dtype) { case DType::kF32: AttentionDenseFlashKernel<<>>( @@ -3376,9 +3418,15 @@ void AttentionDenseFlashKernelCuda(Queue& q, Tensor& out, const Tensor& query, c // The fast path is deliberately NARROW — bf16, head_dim 64, non-causal, MHA, and the // vendored kernels compiled in — because head_dim 64 non-split is the only extra // instantiation this change compiles. Anything else falls through to -// AttentionDenseFlash, so callers get the best available kernel for their shape rather -// than a hard refusal, and every non-encoder caller of this op is byte-identical to +// AttentionDenseFlash, and every non-encoder caller of this op is byte-identical to // the flash-tiled path by construction. +// +// The fall-through is not a promise that every shape runs. AttentionDenseFlash has +// its own head_dim domain — the K/V tile's shared-memory request against CUDA's +// default 48 KiB cap, so 192 in bf16 and 96 in f32 (#1544) — and refuses above it +// naming vt::AttentionDenseFast. A shape wider than that reaches a NAMED refusal +// through here, not a kernel. Every caller of this op today is far inside the bound +// (max head_dim 80), so nothing that runs takes that path. #ifdef VLLM_CPP_FLASH_ATTN // Same-binary A/B + RED knob, mirroring VT_FA2_PREFILL / VT_FA2_DECODE // (cuda_paged_attn.cu): VT_FA2_DENSE=0 restores the scalar flash-tiled kernel so both diff --git a/tests/scripts/test_check_attention_rung_consistency.py b/tests/scripts/test_check_attention_rung_consistency.py new file mode 100755 index 000000000..c4179d934 --- /dev/null +++ b/tests/scripts/test_check_attention_rung_consistency.py @@ -0,0 +1,438 @@ +#!/usr/bin/env python3 +"""Unit and mutation checks for scripts/check-attention-rung-consistency.py. + +The mutation cases below are the point of the file. A checker that reports zero +drift on a green tree proves nothing on its own: it reports zero drift when its +regex matches nothing at all, which is exactly how #1544's defect went unseen for +nine call sites. `MutationTests` below holds the cases that guard against it. +""" + +from __future__ import annotations + +import contextlib +import importlib.util +import io +import sys +import tempfile +import unittest +from pathlib import Path +from unittest import mock + + +ROOT = Path(__file__).resolve().parents[2] +CHECKER = ROOT / "scripts/check-attention-rung-consistency.py" +SPEC = importlib.util.spec_from_file_location("check_attention_rung", CHECKER) +assert SPEC is not None and SPEC.loader is not None +mod = importlib.util.module_from_spec(SPEC) +sys.modules[SPEC.name] = mod +SPEC.loader.exec_module(mod) + +MODELS = "src/vllm/model_executor/models" +ALLOWLIST = ROOT / "scripts/attention-rung-allowlist.txt" + +MARKED = """\ +void Forward() { + // VT-ATTN-NAIVE: reference arm of the dense/paged equivalence gate. + vt::Attention(q, out, qq, kk, vv, args); +} +""" + +UNMARKED = """\ +void Forward() { + vt::Attention(q, out, qq, kk, vv, args); +} +""" + + +class CallDetectionTests(unittest.TestCase): + def test_unmarked_site_is_found(self) -> None: + self.assertEqual(mod.scan_file(UNMARKED), [(2, False)]) + + def test_marked_site_is_found_and_credited(self) -> None: + self.assertEqual(mod.scan_file(MARKED), [(3, True)]) + + def test_fast_rungs_are_never_sites(self) -> None: + # The trailing `\(` is what excludes these, not the `\b`. Measured: all + # five still fail to match with the `\b` removed. + for fast in ( + "vt::AttentionDenseFlash(q, o, a, b, c, args);", + "vt::AttentionDenseFast(q, o, a, b, c, args);", + "vt::AttentionDenseFa2(q, o, a, b, c, args);", + "vt::AttentionCross(q, o, a, b, c, args);", + "vt::PagedAttention(q, o, a, b, c, args);", + ): + self.assertEqual(mod.scan_file(fast), [], fast) + + def test_a_commented_out_call_is_not_a_site(self) -> None: + self.assertEqual(mod.scan_file("// vt::Attention(q, o, a, b, c, args);\n"), []) + + def test_a_block_commented_call_is_not_a_site(self) -> None: + text = "/* legacy:\n vt::Attention(q, o, a, b, c, args);\n*/\n" + self.assertEqual(mod.scan_file(text), []) + + def test_an_if_zero_call_is_not_a_site(self) -> None: + text = "#if 0\nvt::Attention(q, o, a, b, c, args);\n#endif\n" + self.assertEqual(mod.scan_file(text), []) + + def test_line_numbers_survive_normalization(self) -> None: + # normalize_source is position-preserving; a checker that reported a line + # from the normalized text would drift by every stripped block comment. + text = "/* a\n b\n c */\n\nvt::Attention(q, o, a, b, c, args);\n" + self.assertEqual(mod.scan_file(text), [(5, False)]) + self.assertEqual(text.splitlines()[4].strip()[:14], "vt::Attention(") + + def test_two_sites_are_reported_independently(self) -> None: + text = MARKED + "\n" * 40 + UNMARKED + sites = mod.scan_file(text) + self.assertEqual([marked for _, marked in sites], [True, False]) + + +class MarkerTests(unittest.TestCase): + def test_reason_must_be_substantive(self) -> None: + self.assertIsNone(mod.marker_reason(" // nothing to see here")) + self.assertEqual(mod.marker_reason("// VT-ATTN-NAIVE: x"), "x") + # ...but a one-character reason does not satisfy the site. + self.assertFalse(mod.has_marker(["// VT-ATTN-NAIVE: x", "vt::Attention(a);"], 2)) + + def test_marker_must_be_a_comment(self) -> None: + # A string literal naming the marker is not a record. + self.assertIsNone(mod.marker_reason('const char* s = "VT-ATTN-NAIVE: nope at all";')) + + def test_marker_window_is_bounded(self) -> None: + lines = ["// VT-ATTN-NAIVE: a genuine recorded reason"] + [""] * 40 + lines.append("vt::Attention(a);") + self.assertFalse(mod.has_marker(lines, len(lines))) + near = ["// VT-ATTN-NAIVE: a genuine recorded reason"] + [""] * 5 + near.append("vt::Attention(a);") + self.assertTrue(mod.has_marker(near, len(near))) + + def test_marker_on_the_call_line_counts(self) -> None: + line = "vt::Attention(a); // VT-ATTN-NAIVE: the eager rung of the A/B" + self.assertTrue(mod.has_marker([line], 1)) + + +class DriftTests(unittest.TestCase): + def test_marked_site_never_drifts(self) -> None: + self.assertEqual( + mod.drift_sites({f"{MODELS}/nemotron_h.cpp": [(675, True)]}, set()), [] + ) + + def test_unmarked_site_drifts(self) -> None: + self.assertEqual( + mod.drift_sites({f"{MODELS}/muse_glimmer_vision.cpp": [(639, False)]}, set()), + [(f"{MODELS}/muse_glimmer_vision.cpp", 639)], + ) + + def test_allowlisted_stem_passes(self) -> None: + self.assertEqual( + mod.drift_sites({f"{MODELS}/ltx2.cpp": [(959, False)]}, {"ltx2"}), + [], + ) + + def test_mixed_reports_only_the_unmarked(self) -> None: + self.assertEqual( + mod.drift_sites( + { + f"{MODELS}/whisper_audio.cpp": [(324, True)], + f"{MODELS}/qwen3_5.cpp": [(5279, True)], + f"{MODELS}/muse_glimmer_vision.cpp": [(639, False)], + f"{MODELS}/ltx2.cpp": [(959, False)], + }, + allowlisted={"ltx2"}, + ), + [(f"{MODELS}/muse_glimmer_vision.cpp", 639)], + ) + + def test_stale_entry_is_reported_and_not_fatal(self) -> None: + scanned = {f"{MODELS}/ltx2.cpp": [(959, True)]} + self.assertEqual(mod.stale_allowlist_entries(scanned, {"ltx2"}), ["ltx2"]) + self.assertEqual(mod.drift_sites(scanned, {"ltx2"}), []) + + def test_a_header_and_a_cpp_sharing_a_stem_do_not_collide(self) -> None: + # Keyed on the PATH: keyed on the stem, ltx2.h would overwrite ltx2.cpp and + # the checker would silently scan one file instead of two. + scanned = { + f"{MODELS}/ltx2.cpp": [(959, False)], + "include/vllm/model_executor/models/ltx2.h": [(31, False)], + } + self.assertEqual(len(scanned), 2) + self.assertEqual(len(mod.drift_sites(scanned, set())), 2) + self.assertEqual(mod.drift_sites(scanned, {"ltx2"}), []) + + def test_allowlist_parsing(self) -> None: + text = "# comment\nltx2 # trailing reason\nltx2_device\n\n" + self.assertEqual(mod.allowlisted_names(text), {"ltx2", "ltx2_device"}) + + +class ShippedTreeTests(unittest.TestCase): + def scan(self): + return mod.scan_models(), mod.allowlisted_names( + ALLOWLIST.read_text(encoding="utf-8") + ) + + def test_shipped_tree_is_green(self) -> None: + scanned, allowed = self.scan() + self.assertEqual(mod.drift_sites(scanned, allowed), []) + + def test_the_population_is_not_empty(self) -> None: + # A scanner that matches nothing is green for the wrong reason: an empty + # scan and a clean tree file the same report. This is the guard against a + # regex that stops matching after a rename, and it asserts only that the + # scan still finds SOMETHING. + # + # It deliberately does NOT pin the count. A raw total is a measurement of + # the model tree stored in this file, so it reds on every row that + # legitimately REMOVES a vt::Attention call -- which is precisely the rows + # the allowlist exists to unblock, and precisely the drift lock AGENTS.md + # `## Records` forbids: never store a measurement of one file inside + # another file. The floor of 9 this replaces had zero headroom against a + # shipped tree of exactly 9 sites, so #1545's routing change alone would + # have turned `main` red. Issue #1629. + # + # What a count would have bought is covered without the coupling: + # test_the_six_deliberate_sites_carry_a_marker names its files, and + # test_every_allowlisted_stem_names_a_real_model_source below catches the + # bogus entry a total never could. + scanned, _ = self.scan() + self.assertGreaterEqual(sum(len(v) for v in scanned.values()), 1) + + def test_the_six_deliberate_sites_carry_a_marker(self) -> None: + scanned, _ = self.scan() + for stem in ( + "whisper_audio", + "qwen3_vl_vision", + "kimi_linear_device", + "qwen3_5", + "nemotron_h", + "nemotron_h_device", + ): + path = f"{MODELS}/{stem}.cpp" + self.assertIn(path, scanned, path) + self.assertTrue(all(m for _, m in scanned[path]), path) + + def test_allowlist_holds_only_the_in_flight_stems(self) -> None: + # It is not a parking lot. Growth is a review decision, and this pins the + # set so growth is visible in a diff of this file. + _, allowed = self.scan() + self.assertEqual(allowed, {"muse_glimmer_vision", "ltx2", "ltx2_device"}) + + def test_every_allowlisted_stem_names_a_real_model_source(self) -> None: + # Pins that every allowlisted stem names a model source that exists, so a + # typo is reported at the typo. + # + # Keyed on the FILE existing, never on scan membership. A stem stops having + # a call site the moment its removing row lands -- that is the state the + # allowlist is built to survive, which the checker's own + # stale_allowlist_entries docstring states -- so asserting the stem is + # still in `scanned` would rebuild exactly the lock #1629 removes. + _, allowed = self.scan() + for stem in sorted(allowed): + sources = [ + models_dir / f"{stem}{suffix}" + for models_dir in mod.MODEL_DIRS + for suffix in (".cpp", ".h") + ] + self.assertTrue( + any(path.is_file() for path in sources), + f"allowlisted stem {stem!r} names no model source: none of " + + ", ".join(str(path) for path in sources) + + " exists. It excuses nothing, and the checker reports the " + "mismatch only indirectly: an unexcused call site in the file " + "the stem was meant to cover, or a STALE line naming the " + "misspelling.", + ) + + +class MutationTests(unittest.TestCase): + """Guards for the regressions the checker exists to catch.""" + + def setUp(self) -> None: + self.scanned, self.allowed = mod.scan_models(), mod.allowlisted_names( + ALLOWLIST.read_text(encoding="utf-8") + ) + + def test_a_new_unmarked_model_goes_red(self) -> None: + mutated = dict(self.scanned) + new = f"{MODELS}/some_new_vision_tower.cpp" + mutated[new] = [(412, False)] + self.assertEqual(mod.drift_sites(mutated, self.allowed), [(new, 412)]) + + def test_a_new_unmarked_call_in_a_HEADER_goes_red(self) -> None: + # The bypass the .h glob closes: a call moved into an inline function. + mutated = dict(self.scanned) + hdr = "include/vllm/model_executor/models/some_new_tower.h" + mutated[hdr] = [(88, False)] + self.assertEqual(mod.drift_sites(mutated, self.allowed), [(hdr, 88)]) + + def test_deleting_a_marker_goes_red(self) -> None: + mutated = dict(self.scanned) + path = f"{MODELS}/whisper_audio.cpp" + mutated[path] = [(line, False) for line, _ in self.scanned[path]] + self.assertTrue(mod.drift_sites(mutated, self.allowed)) + + def test_a_second_unmarked_call_in_a_marked_file_goes_red(self) -> None: + # Per-SITE, not per-file: a file that already records one reason must not + # launder a new naive call added elsewhere in it. + mutated = dict(self.scanned) + path = f"{MODELS}/qwen3_5.cpp" + mutated[path] = list(self.scanned[path]) + [(9999, False)] + self.assertIn((path, 9999), mod.drift_sites(mutated, self.allowed)) + + def test_a_stub_reason_goes_red(self) -> None: + lines = ["// VT-ATTN-NAIVE: todo", "vt::Attention(a);"] + self.assertFalse(mod.has_marker(lines, 2)) + + def test_widening_the_regex_to_the_fast_rungs_is_visible(self) -> None: + # These pin the trailing `\(` and the `\s*`, never the `\b`. Measured: with + # the `\b` removed, this suite and the shipped tree both stay green. + # + # Three checker comments name a cause that measurement refutes, and none is + # repairable here: check-pr-size.py refuses a comment-only edit to a + # governance checker (#1631). They are the widening paragraph (:58-61), the + # `_NAIVE_CALL` comment (:93-96), and the `excused` cause (:252-255). + self.assertIsNone(mod._NAIVE_CALL.search("vt::AttentionDenseFlash(a);")) + self.assertIsNotNone(mod._NAIVE_CALL.search("vt::Attention (a);")) + + +class GreenReportTests(unittest.TestCase): + """What the OK line tells a reader who never opens the allowlist. + + A green that prints only "9 sites, 6 marked" reads as three unaccounted sites + or as nothing at all, depending on whether the reader does the subtraction. The + number that matters is how many sites carry NO reason and pass anyway. Nothing + asserted this line before, so the count could be dropped or silently go wrong + without a red. + """ + + def report(self) -> str: + buffer = io.StringIO() + with contextlib.redirect_stdout(buffer): + code = mod.main() + self.assertEqual(code, 0, buffer.getvalue()) + return buffer.getvalue() + + def report_over(self, scanned, allowlist_text: str) -> str: + """`main()`'s own report path, driven over a CONSTRUCTED tree. + + The report reads two module-level names, so both are redirected for the + call: `scan_models`, to a dict built by hand exactly as `MutationTests` + builds one, and `ALLOWLIST`, to a temporary file. Nothing the model tree + or the real allowlist does can then change what these cases assert, which + is the whole reason they do not read either one. + """ + buffer = io.StringIO() + with tempfile.TemporaryDirectory() as tmp: + allowlist = Path(tmp) / "attention-rung-allowlist.txt" + allowlist.write_text(allowlist_text, encoding="utf-8") + with mock.patch.object(mod, "scan_models", lambda: scanned): + with mock.patch.object(mod, "ALLOWLIST", allowlist): + with contextlib.redirect_stdout(buffer): + code = mod.main() + self.assertEqual(code, 0, buffer.getvalue()) + return buffer.getvalue() + + def test_the_ok_line_counts_the_sites_an_allowlist_excuses(self) -> None: + # The guard proper: the "unmarked and excused" branch is exercised with a + # NON-ZERO count on a tree this file owns. The constructed scan holds one + # allowlisted file carrying an unmarked site -- the debt the green hides -- + # beside a marked one, plus a marked site nothing excuses, so a checker + # that dropped the clause, or hard-coded the number, cannot produce this + # line. Both mutations were run against this case and both turn it red. + # + # A checker that printed `sites - marked` in place of `excused` is NOT on + # that list, and no constructed scan can add it. `main()` reaches the OK + # line only when `drift_sites` is empty, and `drift_sites` is empty exactly + # when no unmarked site sits outside an allowlisted file -- so on every + # green the checker can print, every unmarked site is excused and + # `excused == sites - marked` identically. That substitution was mutated + # into the checker and the whole suite stayed green, so the claim is + # recorded here as unpinnable rather than left standing as a guarantee this + # case does not carry. + report = self.report_over( + { + f"{MODELS}/ltx2.cpp": [(10, False), (20, True)], + f"{MODELS}/whisper_audio.cpp": [(30, True)], + }, + "# in flight\nltx2 # a row already open removes this call\n", + ) + self.assertEqual( + report.strip(), + "OK (attention rung): 3 vt::Attention call site(s) in 2 model source " + "file(s); 2 carry a recorded reason, 1 unmarked and excused by 1 " + "allowlisted in-flight stem(s).", + ) + + def test_the_ok_line_reports_zero_when_no_stem_is_allowlisted(self) -> None: + # The state the allowlist exists to REACH: every in-flight row has landed + # and the file parks no stem. The checker prints the count even at zero, + # and this pins that, so no case has to require the shipped tree to still + # carry excused debt in order to keep the branch covered. + report = self.report_over( + {f"{MODELS}/whisper_audio.cpp": [(30, True)]}, + "# nothing in flight\n", + ) + self.assertEqual( + report.strip(), + "OK (attention rung): 1 vt::Attention call site(s) in 1 model source " + "file(s); 1 carry a recorded reason, 0 unmarked and excused by 0 " + "allowlisted in-flight stem(s).", + ) + + def test_the_ok_line_reports_the_excused_sites(self) -> None: + # The shipped-tree half: whatever the tree's excused count IS, the OK line + # must state it. `excused` is RE-DERIVED here and never pinned, so this + # holds at 3 today and at 0 once the last in-flight stem is cleaned up. + # + # It deliberately does NOT assert `excused > 0`. That floor was a + # measurement of the model tree stored in this file: it passes only while + # some stem is still parked on the allowlist, so the row that removes the + # LAST one turns this case red while the checker itself is green at rc=0 -- + # the drift lock AGENTS.md `## Records` forbids, and the same shape as the + # `>= 9` population floor #1629 removed. The branch it was standing in for + # is covered above, on a tree this file constructs. + scanned = mod.scan_models() + allowed = mod.allowlisted_names(ALLOWLIST.read_text(encoding="utf-8")) + excused = sum( + 1 + for path, sites in scanned.items() + if Path(path).stem in allowed + for _, marked in sites + if not marked + ) + self.assertIn(f"{excused} unmarked and excused by", self.report()) + + def test_the_excused_count_is_not_sites_minus_marked(self) -> None: + # The subtraction a reader would do by hand. + # + # Read this as a record of the two DEFINITIONS, not as a gate on the OK + # line: it never calls `main()`, and it derives both numbers itself. + # Holding the OK line to `excused` rather than to the subtraction is not + # something this file can do from any scan -- the identity is spelled out + # in test_the_ok_line_counts_the_sites_an_allowlist_excuses. + scanned = { + "src/vllm/model_executor/models/ltx2.cpp": [(10, False), (20, True)], + "src/vllm/model_executor/models/whisper_audio.cpp": [(30, True)], + } + allowed = {"ltx2"} + excused = sum( + 1 + for path, sites in scanned.items() + if Path(path).stem in allowed + for _, marked in sites + if not marked + ) + sites = sum(len(v) for v in scanned.values()) + marked = sum(1 for v in scanned.values() for _, m in v if m) + self.assertEqual(excused, 1) + self.assertEqual(sites - marked, 1) + # They agree HERE, and on any scan the checker reports OK on, because no + # unmarked site sits outside the allowlist -- which is what the assert + # below states. Dropping the allowlisted file's marked site does not + # separate them either: it lowers `sites` and `marked` together. They + # diverge only once an unmarked site lands in a file no stem excuses, and + # that scan is a red. + self.assertEqual(mod.drift_sites(scanned, allowed), []) + + +if __name__ == "__main__": + unittest.main(verbosity=2) diff --git a/tests/scripts/test_check_pr_size.py b/tests/scripts/test_check_pr_size.py index 201c092c0..a19dd74ae 100755 --- a/tests/scripts/test_check_pr_size.py +++ b/tests/scripts/test_check_pr_size.py @@ -507,6 +507,14 @@ def test_every_created_checker_has_closed_bootstrap_evidence(self) -> None: # them; every case that reads an exit code of 1 or an examined count # goes red. "scripts/check-conflict-markers.py", + # 2026-08-21: the attention-rung gate (#1544). Created in the same + # range, so it has no BASE version to mutate. Its suite loads the + # checker as a module at import time and every case then calls into + # it, so the disabled stub -- which defines none of scan_file, + # has_marker, drift_sites, stale_allowlist_entries or main -- takes + # all 31 cases red on AttributeError. Measured, not asserted: the + # suite has no case that passes without touching the checker. + "scripts/check-attention-rung-consistency.py", } self.assertEqual(set(checker.CREATION_MUTATIONS), expected) for path, mutation in checker.CREATION_MUTATIONS.items(): diff --git a/tests/vt/test_ops_attention.cpp b/tests/vt/test_ops_attention.cpp index 5dffc05a4..6cea50739 100644 --- a/tests/vt/test_ops_attention.cpp +++ b/tests/vt/test_ops_attention.cpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include "vt/backend.h" @@ -349,3 +350,99 @@ TEST_CASE("attention CUDA matches CPU (causal, GQA, real head_dim)") { RunCudaCase(/*T=*/6, /*Hq=*/4, /*Hk=*/2, /*D=*/8, 0.35f, true, 77); RunCudaCase(/*T=*/5, /*Hq=*/2, /*Hk=*/1, /*D=*/16, 0.25f, /*causal=*/false, 999); } + +// --- AttentionDenseFlash's head_dim contract (#1544) ------------------------- +// +// The op advertised `head_dim <= 256` while requesting `2*64*d*sizeof(Tin)` bytes +// of dynamic shared memory with no `cudaFuncSetAttribute` anywhere in +// src/vt/cuda/, so the driver's default 48 KiB cap made the real ceiling 192 in +// bf16 and 96 in f32. A caller at head_dim 256 got a bare launch failure naming +// nothing it could do instead. These cases pin the arithmetic that fixes it, and +// they run with no GPU because that arithmetic is a pure host function — which is +// the whole reason it lives in include/vt/ops.h rather than inside the .cu. +TEST_CASE("attention-dense-flash: the advertised head_dim is the launchable one") { + // The tile the CTA stages: K and V, 64 columns each, head_dim wide. + CHECK(vt::AttentionDenseFlashSmemBytes(64, 2) == 16384); + CHECK(vt::AttentionDenseFlashSmemBytes(96, 2) == 24576); // Muse Glimmer, #1545 + CHECK(vt::AttentionDenseFlashSmemBytes(192, 2) == 49152); // exactly the cap + CHECK(vt::AttentionDenseFlashSmemBytes(256, 2) == 65536); // over it + CHECK(vt::AttentionDenseFlashSmemBytes(96, 4) == 49152); // exactly the cap + CHECK(vt::AttentionDenseFlashSmemBytes(192, 4) == 98304); // Kimi, over it + CHECK(vt::AttentionDenseFlashSmemBytes(256, 4) == 131072); + + // The honest bounds. These are the widths the launcher now refuses above. + CHECK(vt::AttentionDenseFlashMaxHeadDim(/*bf16=*/2) == 192); + CHECK(vt::AttentionDenseFlashMaxHeadDim(/*f32=*/4) == 96); + + // Neither is the 256 the op used to advertise, and that gap IS the defect. + CHECK(vt::AttentionDenseFlashMaxHeadDim(2) < vt::kAttentionDenseMaxHeadDim); + CHECK(vt::AttentionDenseFlashMaxHeadDim(4) < vt::kAttentionDenseMaxHeadDim); + + // The bound is INCLUSIVE at the cap and exclusive one element past it. head_dim + // 192 in bf16 sits exactly on 49152 and launches today, so a bound that refused + // it would be a regression rather than a repair. + CHECK(vt::AttentionDenseFlashSmemBytes(vt::AttentionDenseFlashMaxHeadDim(2), 2) <= + vt::kCudaDefaultDynamicSmemBytes); + CHECK(vt::AttentionDenseFlashSmemBytes(vt::AttentionDenseFlashMaxHeadDim(2) + 1, 2) > + vt::kCudaDefaultDynamicSmemBytes); + CHECK(vt::AttentionDenseFlashSmemBytes(vt::AttentionDenseFlashMaxHeadDim(4), 4) <= + vt::kCudaDefaultDynamicSmemBytes); + CHECK(vt::AttentionDenseFlashSmemBytes(vt::AttentionDenseFlashMaxHeadDim(4) + 1, 4) > + vt::kCudaDefaultDynamicSmemBytes); + + // The register blocking still bounds it from the other side: a hypothetical + // 1-byte input would clear the shared-memory cap far past what the kernel can + // hold in registers, and the bound must report the register limit there. + CHECK(vt::AttentionDenseFlashMaxHeadDim(1) == vt::kAttentionDenseMaxHeadDim); +} + +TEST_CASE("attention-dense-flash: an over-cap head_dim is REFUSED, naming the rung") { + // This one needs a device: the refusal lives in the CUDA launcher, and nothing + // on a CPU-only box executes it. Loud on purpose — a quiet skip here would let + // the launcher lose the bound with every gate still green (spec R1, #1573). + if (!HasCuda()) { + MESSAGE("no CUDA backend; the AttentionDenseFlash head_dim refusal is PENDING"); + return; + } + const int64_t T = 4, Hq = 2, Hk = 2, D = 256; // f32 => a 128 KiB K/V tile + auto q = RandF32(static_cast(T * Hq * D), 5150); + auto k = RandF32(static_cast(T * Hk * D), 5151); + auto v = RandF32(static_cast(T * Hk * D), 5152); + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + QueueGuard g(gpu); + DeviceTensor dq(gpu, g.q, DType::kF32, {T, Hq, D}, q.data()); + DeviceTensor dk(gpu, g.q, DType::kF32, {T, Hk, D}, k.data()); + DeviceTensor dv(gpu, g.q, DType::kF32, {T, Hk, D}, v.data()); + DeviceTensor dout(gpu, g.q, DType::kF32, {T, Hq, D}); + bool threw = false; + std::string what; + try { + vt::AttentionDenseFlash(g.q, dout.tensor(), dq.tensor(), dk.tensor(), dv.tensor(), + AttentionArgs{std::pow(256.0f, -0.5f), /*causal=*/false}); + } catch (const std::runtime_error& e) { + threw = true; + what = e.what(); + } + CHECK(threw); + // Refusing is only half of it. The message must name the rung that DOES serve + // this width, or the caller is left where the opaque launch error left them. + CHECK(what.find("AttentionDenseFast") != std::string::npos); + CHECK(what.find("head_dim") != std::string::npos); + + // ...and the honest bound is not a blanket refusal: head_dim 96 in f32 lands + // exactly on the cap and must still run. + const int64_t Dok = 96; + auto q2 = RandF32(static_cast(T * Hq * Dok), 5153); + auto k2 = RandF32(static_cast(T * Hk * Dok), 5154); + auto v2 = RandF32(static_cast(T * Hk * Dok), 5155); + DeviceTensor dq2(gpu, g.q, DType::kF32, {T, Hq, Dok}, q2.data()); + DeviceTensor dk2(gpu, g.q, DType::kF32, {T, Hk, Dok}, k2.data()); + DeviceTensor dv2(gpu, g.q, DType::kF32, {T, Hk, Dok}, v2.data()); + DeviceTensor dout2(gpu, g.q, DType::kF32, {T, Hq, Dok}); + vt::AttentionDenseFlash(g.q, dout2.tensor(), dq2.tensor(), dk2.tensor(), + dv2.tensor(), + AttentionArgs{std::pow(96.0f, -0.5f), /*causal=*/false}); + std::vector got(static_cast(T * Hq * Dok), 0.0f); + dout2.Download(g.q, got.data()); + CHECK(std::isfinite(got[0])); +}