fix(KERNEL-ATTN-DENSE-FLASH): a model on the naive attention kernel now says why, and AttentionDenseFlash advertises the head_dim it can launch (#1544)

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

.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

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

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

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: <why this call must stay on kAttention>
+#
+# 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/<stem>.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.