fix(rocm): #785 runtime arch dispatch — launch SharedKWmma prefill where it exists

.agents/issue-index.md

@@ -380,3 +380,4 @@ rather than merged. `scripts/check-agent-record.py` gates both.
 | [#1225](https://github.com/mudler/vllm.cpp/issues/1225) | `SPEC-DSPARK-BLOCK-SIZE-GUARD` | A DSpark speculative length below the draft's block is accepted silently. `ResolveDspark` carries upstream's `k >= dspark_block_size` hard error (`include/vllm/config/speculative.h:179-185`, from `vllm/config/speculative.py:1003-1027` @ `555967922`) and both production call sites pass `std::nullopt` for `n_predict` and for `dspark_block_size` (`src/vllm/entrypoints/model_loader.cpp:881-883` and `:1675-1677`), so the floor reaches no user path and only `tests/vllm/config/test_speculative_dspark.cpp:99-107` drives it. Nothing in our draft path reads the checkpoint's block key — the block layout is sized by `k` alone (`include/vllm/v1/worker/gpu/spec_decode/dspark/speculator.h:56`) and no weight is block-shaped — so a short `k` raises no shape error and drafts a structurally wrong block while the tokens keep flowing. A literal port does NOT close it: `dspark_block_size` appears in no pinned file but `speculative.py`, and neither published Qwen3 draft sets it (`deepseek-ai/dspark_qwen3_4b_block7` and `RadixArk/Qwen3.8-27B-DSpark` @ `85ef153b` both carry `block_size: 7`, no `n_predict`), while upstream's `block_size` normalization at `:945-961` is Gemma4-only — so upstream accepts `k=6` on both sides of vllm#52197. Closing it for the lane we ship needs `block_size` as the floor fallback, one tracked divergence argued in the spec and the commit | bug |
 | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | `ENG-CUDAGRAPH-DEDUP` | The owed device A/B for #1162 RAN on `dgx:gpu0` 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f`) and it SPLIT. CORRECTNESS PASSED and closes [#1184](https://github.com/mudler/vllm.cpp/issues/1184): 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and three workloads hashing to three DIFFERENT values. THE BENEFIT IS REFUTED for the case the row was filed for: `N == M` in every `VT_CUDA_GRAPH_DEDUP=1` cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact of the first attempt. The cause is structural and was pre-registered before the run: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension is IN THE KEY, no candidate group forms and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own premise, and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded — clocks unpinned AND the ON arm allocated exactly as many executables as OFF. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF and the default flip is NOT JUSTIFIED on this evidence. THE OPEN HYPOTHESIS this issue owns, deliberately not decided by the record that filed it: a COARSER key keeping the function addresses and topology but dropping the launch dimensions and memcpy extents would let two padded buckets form a candidate group at all, and the probe-before-fold design means a wrong grouping costs one wasted `cudaGraphExecUpdate` probe and a private executable rather than a wrong replay, so it is a cost question needing its own spec, a red-first discrimination test and a device probe-refusal rate. Also owed here: whether ANY other capture site or two models sharing the process-singleton registry can fold (only the Qwen3 dense decode driver was exercised), and a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`; a CUDA 12.x toolkit is the untried route). Honest gaps carried into the record: per-shape replay counts are unavailable, so B's ~30-per-shape is arithmetic over a printed TOTAL; the driver's "N captured size(s)" counts SLOTS not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0. Full evidence [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W4`; spec [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` | perf |
 | [#1229](https://github.com/mudler/vllm.cpp/issues/1229) | `ENG-CUDAGRAPH-DEDUP` | `origin/main` went RED at `fe24a3029` because `tests/scripts/test_check_gate_commands.py:670` pinned the literal string `Device byte-identity A/B (owed` inside `.agents/specs/eng-cudagraph-dedup.md`, and that commit legitimately rewrote the sentence from recording the device A/B as owed to recording that it RAN and what it found. The record edit is correct and the assertion described a state the record left, so every branch that merged `origin/main` inherited the failure. The protected property is stated by the assertion's own comment: the row is credited in `RUNNABLE_BASELINE` for `ctest -R test_graph_dedup` and `./scripts/agent-preflight.sh`, both CPU-tier, while gate item 6 of the spec contributes NO runnable command, so the credit covers the CPU tier and the record must not go silent about the leg it does not cover. That property is STATE-INDEPENDENT: `owed` was the leg's value on the day, not the property, and silence is the defect whether the leg was never run or the line was deleted after it ran. Asserting the NEW wording would only move the pin from `(owed` to `RAN 2026-08-18, PASS`, and the rerun this row already owes under [#1226](https://github.com/mudler/vllm.cpp/issues/1226) against a coarser signature key would red it again. FIXED by keying on structure: `gate_items`, `item_lead` and `gate_disposition` land in `scripts/check-gate-commands.py`, and the test now asserts that exactly one gate item names the device byte-identity A/B as its SUBJECT, that `runnable_commands` extracts nothing from it, and that its bold lead declares a disposition from a closed vocabulary. Lead-scoping is measured rather than assumed: a whole-item search credits item 5 with a disposition it does not declare, because its body contains `fail` in the phrase "proving the instrument can fail". Both historical wordings of item 6 are pinned as fixtures, which is what proves the new assertion is state-independent rather than re-pinned to today's sentence. The rule is NOT swept over every gate item and the reason is measured, not assumed: items 1 and 5 yield no command and declare no disposition, so a sweep is red on arrival, and widening the vocabulary until they pass would make it match ordinary English. That general rule is carried under `## Owed`. Also established and NOT a defect: `ERROR: these baseline rows left the gated population: ROW-THAT-IS-NOT-THERE` on the same run is the expected negative-path output of `test_check_mode_is_never_silently_swallowed_by_json`, confirmed by running that fixture alone. Spec [`fix-gate-commands-prose-pin.md`](specs/fix-gate-commands-prose-pin.md) | bug |
+| [#785](https://github.com/mudler/vllm.cpp/issues/785) | `BACKEND-ROCM` | ROCm: `PagedAttnPrefillSharedKWmma` host launches were behind `#if defined(VT_ROCWMMA_OK)`, which is never true on the host pass. This row repairs d=256 host dispatch only (`<2,8,16,32,false>`). d=512 stays scalar SharedK because shipping-f58b WMMA spills (52/212) and is a separate kernel repair. Spec [`rocm-sharedk-wmma-host-dispatch.md`](specs/rocm-sharedk-wmma-host-dispatch.md) | bug |

.agents/specs/rocm-sharedk-wmma-host-dispatch.md

@@ -0,0 +1,71 @@
+# ROCm: launch SharedK WMMA from the host (#785)
+
+Row: `BACKEND-ROCM`. Issue:
+[#785](https://github.com/mudler/vllm.cpp/issues/785).
+
+## Defect
+
+`PagedAttnPrefillSharedKWmma` launches were behind
+`#if defined(VT_ROCWMMA_OK)`. That macro is defined only on HIP's
+**device** pass (`__gfx1200__` / `__gfx1201__`). The host pass never
+defines it, so both launch sites were deleted. Every d=256/d=512
+prefill silently ran scalar `PagedAttnPrefillSharedK`.
+
+## P0 (this head)
+
+Repairs **d=256 production dispatch only**. d=512 remains scalar.
+
+- Host launch uses `hipDeviceProp_t.gcnArchName` prefix-match
+  `gfx1200` / `gfx1201` (not substring). `gfx1201:xnack-` matches;
+  `foogfx1201` and `gfx12010` do not.
+- Per-device decision cached once (`std::call_once`). Not per-build.
+  Not getenv.
+- Device kernel-body `#if !defined(VT_ROCWMMA_OK)` stub is unchanged.
+- `VT_ATTN_PREFILL_SHAREDK_WMMA` still forces scalar when `=0`.
+- Host path launches only `PagedAttnPrefillSharedKWmma<2,8,16,32,false>`
+  (d=256). The d=512 WMMA launch/stub is removed, not hidden.
+- d=512 keeps the existing scalar `PagedAttnPrefillSharedK` fallthrough.
+  Shipping f58b d=512 WMMA is VGPR 192 / spill 52 / private 212 — that
+  violates the 0/0 compile gate and is owed as a separate kernel repair.
+- Shipping f58b HIP compile of `rocm_paged_attn.hip` (`clang++` roc-7.2.4
+  `f58b06dce1f9`, `--offload-arch=gfx1201`). Device body unchanged.
+
+Measured KD on that object (fields after `.name`):
+
+| Instantiation | path | vgpr | spill | private | LDS |
+|---|---|---|---|---|---|
+| `<2,8,16,32,false>` | d=256 | 151 | **0** | **0** | 4880 |
+
+P1 (this package; GPU HOLD):
+
+Witness the product `vt::PagedAttention` seam, not a direct WMMA kernel call.
+
+- Fixture: BF16 Q/K/V/out, d=256, qg=2 (hq=2,hk=1), T=64, one request,
+  causal, sliding window left=32 right=0, scale=1/sqrt(256),
+  seeds 78525601/02/03. Tensor SHA-256 frozen in
+  `tests/vt/test_ops_paged_attn_sharedk_wmma_p1.cpp`.
+- A/B same binary, separate processes (`VT_ATTN_PREFILL_SHAREDK_WMMA` is
+  process-static). A = default/on, exact kernel
+  `PagedAttnPrefillSharedKWmma<2,8,16,32,false>`. B = `=0`, that kernel
+  absent, scalar `PagedAttnPrefillSharedK<2,8,...>` present.
+- Both vs the same host f32 oracle. Preregistered BF16 bar:
+  `abs(got-ref) <= 1.5e-2 + 1.0e-2*|ref|`, `corr>=0.999`, no nonfinite.
+  Candidate-vs-scalar distance is reported; bit identity is not required.
+- Fail closed on skip, missing trace, wrong kernel identity, nonfinite,
+  oracle miss, or non-zero status. No silent retry. No timing. No d=512.
+- Runner: `tests/scripts/run-785-p1.sh` (exit 78 without `VT_785_P1_GPU_GO=1`).
+  Trace: `rocprofv3 --kernel-trace` parsed by `tests/scripts/parse_785_p1_trace.py`.
+  Classifier separates family from exact specialization: A = exact
+  WMMA `<2,8,16,32,false>` and no SharedK family and no other WMMA;
+  B = exact scalar `<2,8,32,32>` and no WMMA family and no other scalar;
+  wrong BM/BN, wrong qg/d, mixed, or none = UNKNOWN. GPU binary is an
+  `add_executable` only (not ordinary CTest). Shared fixture header
+  `tests/vt/sharedk_wmma_p1_fixture.h` is used by host and GPU; GPU prints
+  frozen Q/K/V SHA-256 and the runner fails on mismatch. Arms pin
+  `SHAREDK_WMMA=1|0`, SharedK=1, decode-opt=1, decode-GQA=1, CPU-ref=0.
+
+## Owed
+
+- Researcher review of this P1 package, then GPU GO/HOLD
+- Separate kernel repair before any d=512 WMMA launch
+- PR body (never-ran consequence + expected d=256 uplift)

docs/ENVIRONMENT.md

@@ -260,7 +260,7 @@ on CUDA/CPU builds beyond the documented behavior.
 | `VT_GEMMA4_RESIDENT_NATIVE` | unset | Prefer native FP8 resident packs when set |
 | `VT_ATTN_PREFILL_FLASH_SHAREDK` | on | ROCm Gemma-4 prefill: scoreless SharedK flash path (prod default). `0` forces DecodeGqa baseline for A/B |
 | `VT_ATTN_PREFILL_FLASH_WMMA` | off | Lab-only ROCm flash WMMA QK path; leave off for prod |
-| `VT_ATTN_PREFILL_SHAREDK_WMMA` | on | ROCm SharedK fused with rocWMMA QK on sliding d=256 (and global-Q d=512). `0` forces scalar SharedK |
+| `VT_ATTN_PREFILL_SHAREDK_WMMA` | on | ROCm SharedK fused with rocWMMA QK on sliding d=256 only. d=512 stays scalar SharedK (shipping f58b WMMA spills). `0` forces scalar SharedK |
 | `VT_GEMMA4_PREFILL_PEER_ACT` | on | Prefill MoE: run ExpertGeGLU on the expert GPU and peer activations only (not full weight PeerCopy). `0` restores weight PeerCopy |
 | `VT_GEMMA4_GPU0_HEADROOM_GB` | `12` | GiB kept free on GPU0 when packing resident experts (decode vs long-prefill trade). Lab dual R9700 + 49k KV: `8` survives 16k+ prefill; `6` OOMs ~11k |
 | `VT_GEMMA4_PREFILL_BATCH_MOE` | auto / `1` in lab recipe | `=1` group-by-expert prefill GEMM for `T>=64`; `=0` serial M=1 (slow). Unset = auto |

include/vt/rocm/rocm_arch.h

@@ -70,4 +70,22 @@ constexpr std::optional<std::pair<int, int>> CapabilityFromGcnArch(std::string_v
   return std::pair<int, int>{major, minor};
 }
 
+// Host launch gate for Prefill SharedK WMMA (#785). True only when `gcn_arch`
+// is the literal HIP gcnArchName prefix `gfx1200` or `gfx1201`.
+//
+// Prefix, not substring: `foogfx1201` is false. After the six-char stem the
+// next character must be end-of-string or a non-digit so `gfx1201:xnack-`
+// matches and `gfx12010` does not. CapabilityFromGcnArch(12,0) is too wide
+// (gfx1202..gfx1209).
+constexpr bool GcnArchNameIsGfx12PrefillWmma(std::string_view gcn_arch) {
+  auto prefix_ok = [](std::string_view s, std::string_view stem) {
+    if (s.size() < stem.size()) return false;
+    if (s.substr(0, stem.size()) != stem) return false;
+    if (s.size() == stem.size()) return true;
+    const char c = s[stem.size()];
+    return c < '0' || c > '9';
+  };
+  return prefix_ok(gcn_arch, "gfx1200") || prefix_ok(gcn_arch, "gfx1201");
+}
+
 }  // namespace vt::rocm

src/vt/rocm/rocm_paged_attn.hip

@@ -16,8 +16,11 @@
 
 #include "vt/ops.h"
 #include "vt/backend.h"
+#include "vt/rocm/rocm_arch.h"
 
+#include <array>
 #include <cstdlib>
+#include <mutex>
 #include <vector>
 
 namespace vt::rocm {
@@ -32,6 +35,29 @@ inline void Check(hipError_t err, const char* what) {
   }
 }
 
+// Per-device, once: launch SharedK WMMA only when this device's gcnArchName is
+// gfx1200/gfx1201 (prefix). Not per-build, not getenv. Device kernel body
+// still uses VT_ROCWMMA_OK (#785).
+bool PrefillSharedKWmmaHostOk() {
+  int dev = 0;
+  if (hipGetDevice(&dev) != hipSuccess || dev < 0) return false;
+  struct Slot {
+    std::once_flag once;
+    bool ok = false;
+  };
+  static std::array<Slot, 64> slots;
+  if (static_cast<unsigned>(dev) >= slots.size()) return false;
+  std::call_once(slots[static_cast<unsigned>(dev)].once, [dev] {
+    hipDeviceProp_t prop{};
+    if (hipGetDeviceProperties(&prop, dev) != hipSuccess) {
+      slots[static_cast<unsigned>(dev)].ok = false;
+      return;
+    }
+    slots[static_cast<unsigned>(dev)].ok = GcnArchNameIsGfx12PrefillWmma(prop.gcnArchName);
+  });
+  return slots[static_cast<unsigned>(dev)].ok;
+}
+
 // Optional correctness path: run the CPU paged-attn oracle on host copies.
 // VT_ROCM_ATTN_CPU_REF=1 — isolates whether custom HIP attn is the decode bug.
 bool CpuRefEnabled() {
@@ -1736,7 +1762,8 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const
         return e[0] != '0';
       }();
       static const bool sharedk_wmma = [] {
-        // Default ON: d=256 Q-smem WMMA; d=512 global-Q WMMA (K+V only LDS).
+        // Default ON: d=256 Q-smem WMMA only. d=512 stays scalar SharedK
+        // until a separately gated kernel repair (shipping f58b spills).
         // =0 forces scalar SharedK for A/B.
         const char* e = std::getenv("VT_ATTN_PREFILL_SHAREDK_WMMA");
         if (!e) return true;
@@ -1754,10 +1781,10 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const
                   static_cast<size_t>(2 * bn) * static_cast<size_t>(d));
         };
         const unsigned z_groups = (qg == 8 && d == 512) ? 4u : 1u;
-        // Host must not launch WMMA stubs on non-gfx1200/1201 fatbin slices
-        // (kernel body is #else empty there → silent garbage out).
-#if defined(VT_ROCWMMA_OK)
-        if (sharedk_wmma) {
+        // Host must not launch WMMA stubs on non-gfx1200/1201 devices
+        // (kernel body is empty there → silent garbage out). Runtime prefix
+        // match on hipDeviceProp_t.gcnArchName; VT_ROCWMMA_OK is device-only.
+        if (sharedk_wmma && PrefillSharedKWmmaHostOk()) {
           if (d == 256) {
             // Q16K+K16K+V16K=48KB dyn + s_sc~4KB
             constexpr int BM = 16, BN = 32;
@@ -1774,25 +1801,8 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const
                     args.logits_soft_cap, args.causal, window_left, window_right);
             Check(hipGetLastError(), "paged_attention prefill_sharedk_wmma d256 launch");
             return;
-          } else if (d == 512) {
-            // Global Q; dyn K+V only BN16 → 32KB. Partial last tile scalar QK.
-            constexpr int BM = 16, BN = 16;
-            const unsigned tiles = static_cast<unsigned>((total_q + BM - 1) / BM);
-            dim3 grid(tiles, static_cast<unsigned>(num_kv_heads), z_groups);
-            dim3 block(static_cast<unsigned>(BM * 32));
-            PagedAttnPrefillSharedKWmma<2, 16, BM, BN, true>
-                <<<grid, block, smem_kv(BN), st>>>(
-                    out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(),
-                    k_cache.Ptr<__hip_bfloat16>(), v_cache.Ptr<__hip_bfloat16>(),
-                    block_table.Ptr<int32_t>(), seq_lens.Ptr<int32_t>(),
-                    query_start_loc.Ptr<int32_t>(), num_reqs, hq, num_kv_heads, d, block_size,
-                    bt_row, bt_col, kc_blk, kc_pg, kc_hd, vc_blk, vc_pg, vc_hd, args.scale,
-                    args.logits_soft_cap, args.causal, window_left, window_right);
-            Check(hipGetLastError(), "paged_attention prefill_sharedk_wmma d512 launch");
-            return;
           }
         }
-#endif  // VT_ROCWMMA_OK
         if (d == 256) {
           constexpr int BM = 32, BN = 32;  // K+V 32KB; BM32 ~+2% vs BM16 @11k
           const unsigned tiles = static_cast<unsigned>((total_q + BM - 1) / BM);

tests/CMakeLists.txt

@@ -1530,6 +1530,9 @@ set_tests_properties(test_backend_cross_device_vt_attn_decode_d128 PROPERTIES
 # of that skeleton carrying a DECISION is gated on a CPU-only CI runner with no
 # AMD GPU. The rest of the ROCm skeleton is under VLLM_CPP_HIP below.
 vllm_cpp_add_test(test_rocm_arch vt/test_rocm_arch.cpp)
+# #785 P1 host package (fixture/oracle/trace classifier). CPU-only.
+vllm_cpp_add_test(test_ops_paged_attn_sharedk_wmma_p1
+  vt/test_ops_paged_attn_sharedk_wmma_p1.cpp)
 vllm_cpp_add_test(test_triton_aot_arch_dispatch
   vt/test_triton_aot_arch_dispatch.cpp)
 target_include_directories(test_triton_aot_arch_dispatch PRIVATE
@@ -1539,9 +1542,23 @@ target_include_directories(test_cpu_isa_x86 PRIVATE ${CMAKE_SOURCE_DIR}/src)
 vllm_cpp_add_test(test_cpu_isa_arm vt/test_cpu_isa_arm.cpp)
 target_include_directories(test_cpu_isa_arm PRIVATE ${CMAKE_SOURCE_DIR}/src)
 if(VLLM_CPP_HIP)
-  # Plain C++ (no HIP header): every assertion goes through the vt:: seam. Each
-  # case no-ops when the build has HIP but the box has no AMD GPU.
-  vllm_cpp_add_test(test_rocm_backend vt/test_rocm_backend.cpp)
+# Plain C++ (no HIP header): every assertion goes through the vt:: seam. Each
+# case no-ops when the build has HIP but the box has no AMD GPU.
+vllm_cpp_add_test(test_rocm_backend vt/test_rocm_backend.cpp)
+# #785 P1 GPU product-seam witness. Executable only — NOT add_test.
+# Ordinary CTest must not see this target. Runner fail-closes on 77/nonzero.
+add_executable(test_ops_paged_attn_sharedk_wmma_p1_gpu
+  vt/test_ops_paged_attn_sharedk_wmma_p1_gpu.cpp)
+target_include_directories(test_ops_paged_attn_sharedk_wmma_p1_gpu PRIVATE
+  ${CMAKE_CURRENT_SOURCE_DIR})
+if(TARGET vllm::sanitize_test_shared)
+  target_link_libraries(test_ops_paged_attn_sharedk_wmma_p1_gpu PRIVATE
+    vllm::sanitize_test_shared vllm_test_main)
+else()
+  target_link_libraries(test_ops_paged_attn_sharedk_wmma_p1_gpu PRIVATE
+    vllm::vllm vllm_test_main)
+endif()
+vllm_cpp_set_warnings(test_ops_paged_attn_sharedk_wmma_p1_gpu)
 endif()
 if(VLLM_CPP_METAL)
   vllm_cpp_add_test(test_metal_backend vt/test_metal_backend.cpp)