diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 5c1c465cc..c8ebaf8f3 100644 --- a/.agents/issue-index.md +++ b/.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 | diff --git a/.agents/specs/rocm-sharedk-wmma-host-dispatch.md b/.agents/specs/rocm-sharedk-wmma-host-dispatch.md new file mode 100644 index 000000000..451b68b6c --- /dev/null +++ b/.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) diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index d2fcd9393..2b25ffe00 100644 --- a/docs/ENVIRONMENT.md +++ b/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 | diff --git a/include/vt/rocm/rocm_arch.h b/include/vt/rocm/rocm_arch.h index 4ca94f960..d737bb420 100644 --- a/include/vt/rocm/rocm_arch.h +++ b/include/vt/rocm/rocm_arch.h @@ -70,4 +70,22 @@ constexpr std::optional> CapabilityFromGcnArch(std::string_v return std::pair{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 diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index f661a31da..001ca79f6 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/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 #include +#include #include 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 slots; + if (static_cast(dev) >= slots.size()) return false; + std::call_once(slots[static_cast(dev)].once, [dev] { + hipDeviceProp_t prop{}; + if (hipGetDeviceProperties(&prop, dev) != hipSuccess) { + slots[static_cast(dev)].ok = false; + return; + } + slots[static_cast(dev)].ok = GcnArchNameIsGfx12PrefillWmma(prop.gcnArchName); + }); + return slots[static_cast(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(2 * bn) * static_cast(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((total_q + BM - 1) / BM); - dim3 grid(tiles, static_cast(num_kv_heads), z_groups); - dim3 block(static_cast(BM * 32)); - PagedAttnPrefillSharedKWmma<2, 16, BM, BN, true> - <<>>( - out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), - k_cache.Ptr<__hip_bfloat16>(), v_cache.Ptr<__hip_bfloat16>(), - block_table.Ptr(), seq_lens.Ptr(), - query_start_loc.Ptr(), 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((total_q + BM - 1) / BM); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 863c3dc3f..8777993bf 100644 --- a/tests/CMakeLists.txt +++ b/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) diff --git a/tests/scripts/parse_785_p1_trace.py b/tests/scripts/parse_785_p1_trace.py new file mode 100755 index 000000000..dc2882b6a --- /dev/null +++ b/tests/scripts/parse_785_p1_trace.py @@ -0,0 +1,108 @@ +#!/usr/bin/env python3 +"""#785 P1 kernel-trace classifier. + +Family vs exact specialization (Researcher 2609): + A: exact WMMA <2,8,16,32,false>; no SharedK family; no other WMMA. + B: exact scalar <2,8,32,32>; no WMMA family; no other scalar. + Wrong BM/BN, wrong qg/d, mixed, or none => UNKNOWN. + +WMMA names contain the SharedK prefix; scalar family is scored only on +lines that do not contain SharedKWmma. + +Exit 0 prints `arm=A|B|UNKNOWN` plus marker hits. +Exit 2 = unreadable/empty input (fail closed). +""" +from __future__ import annotations + +import argparse +import sys +from pathlib import Path + +EXACT_WMMA = ( + "PagedAttnPrefillSharedKWmma<2, 8, 16, 32, false>", + "PagedAttnPrefillSharedKWmma<2,8,16,32,false>", + "PagedAttnPrefillSharedKWmmaILi2ELi8ELi16ELi32ELb0E", +) +EXACT_SCALAR = ( + "PagedAttnPrefillSharedK<2, 8, 32, 32>", + "PagedAttnPrefillSharedK<2,8,32,32>", + "PagedAttnPrefillSharedKILi2ELi8ELi32ELi32EE", +) + + +def collect_text(path: Path) -> str: + if path.is_file(): + return path.read_text(errors="replace") + chunks: list[str] = [] + for p in sorted(path.rglob("*")): + if not p.is_file(): + continue + if p.suffix.lower() not in {".csv", ".json", ".txt", ".log", ".out", ""}: + if p.stat().st_size > 8_000_000: + continue + try: + chunks.append(p.read_text(errors="replace")) + except OSError: + continue + return "\n".join(chunks) + + +def _scalar_blob(text: str) -> str: + return "\n".join(line for line in text.splitlines() if "SharedKWmma" not in line) + + +def _strip_exact(text: str, markers: tuple[str, ...]) -> str: + out = text + for m in markers: + out = out.replace(m, "") + return out + + +def classify(text: str) -> tuple[str, list[str], list[str]]: + scalar_blob = _scalar_blob(text) + wmma_hits = [m for m in EXACT_WMMA if m in text] + scalar_hits = [m for m in EXACT_SCALAR if m in scalar_blob] + wmma_family = "PagedAttnPrefillSharedKWmma" in text + scalar_family = "PagedAttnPrefillSharedK" in scalar_blob + other_wmma = "PagedAttnPrefillSharedKWmma" in _strip_exact(text, EXACT_WMMA) + other_scalar = "PagedAttnPrefillSharedK" in _strip_exact(scalar_blob, EXACT_SCALAR) + if wmma_hits and not scalar_family and not other_wmma: + arm = "A" + elif scalar_hits and not wmma_family and not other_scalar: + arm = "B" + else: + arm = "UNKNOWN" + return arm, wmma_hits, scalar_hits + + +def main(argv: list[str]) -> int: + ap = argparse.ArgumentParser(description=__doc__) + ap.add_argument("path", help="rocprofv3 output file or directory") + ap.add_argument("--expect", choices=("A", "B"), help="fail closed if arm mismatches") + args = ap.parse_args(argv) + src = Path(args.path) + if not src.exists(): + print("ERROR: missing trace path", src, file=sys.stderr) + return 2 + text = collect_text(src) + if not text.strip(): + print("ERROR: empty trace", src, file=sys.stderr) + return 2 + arm, wmma_hits, scalar_hits = classify(text) + print(f"arm={arm}") + print("wmma_hits=" + (",".join(wmma_hits) if wmma_hits else "-")) + print("scalar_hits=" + (",".join(scalar_hits) if scalar_hits else "-")) + if args.expect and arm != args.expect: + print(f"ERROR: expected arm={args.expect} got {arm}", file=sys.stderr) + return 1 + if args.expect == "A" and not wmma_hits: + print("ERROR: A missing exact WMMA specialization", file=sys.stderr) + return 1 + if args.expect == "B" and (wmma_hits or not scalar_hits): + print("ERROR: B must be exact scalar-only", file=sys.stderr) + return 1 + return 0 + + +if __name__ == "__main__": + raise SystemExit(main(sys.argv[1:])) diff --git a/tests/scripts/run-785-p1.sh b/tests/scripts/run-785-p1.sh new file mode 100755 index 000000000..4722312b6 --- /dev/null +++ b/tests/scripts/run-785-p1.sh @@ -0,0 +1,111 @@ +#!/usr/bin/env bash +# #785 P1 d=256 SharedK WMMA product-seam witness. +# GPU HOLD until Researcher names GO. Default exit 78. +# +# A: SHAREDK_WMMA=1, exclusive WMMA kernel. +# B: SHAREDK_WMMA=0, exclusive scalar SharedK. +# Same binary, separate processes. No timing. No d=512. Never :8010/:8012. +set -euo pipefail + +if [[ "${VT_785_P1_GPU_GO:-}" != "1" ]]; then + printf '%s\n' "P1 GPU HOLD: set VT_785_P1_GPU_GO=1 only after Researcher GO" + exit 78 +fi + +ROOT="${VT_785_P1_ROOT:-$(cd "$(dirname "$0")/../.." && pwd)}" +BUILD="${VT_785_P1_BUILD:-/home/don/.cache/hermes-builds/vllm-785-p0/build-hip}" +OUT="${VT_785_P1_OUT:-/home/don/.cache/hermes-builds/vllm-785-p0/p1-out}" +BIN="${BUILD}/tests/test_ops_paged_attn_sharedk_wmma_p1_gpu" +PARSER="${ROOT}/tests/scripts/parse_785_p1_trace.py" +ROCPROF="${ROCPROFV3:-/opt/rocm/bin/rocprofv3}" +QHASH=27f164e220edf8d37ddfd1783f0c390968b9d2314ff5d89b7aea3ef50f0eba72 +KHASH=d803e64df022e1cb2049b06d65fc36da02a79ac09f77831e967237e47d49f7ca +VHASH=910548159fd5a5a478573dac6de96f374014cb0b76dac7389071f265dfe1cae7 + +die() { printf 'ERROR: %s\n' "$*" >&2; exit 1; } + +[[ -x "$BIN" ]] || die "missing GPU test binary: $BIN" +[[ -f "$PARSER" ]] || die "missing parser: $PARSER" +[[ -x "$ROCPROF" ]] || die "missing rocprofv3: $ROCPROF" + +rm -rf "$OUT" +mkdir -p "$OUT/A" "$OUT/B" + +run_arm() { + local arm="$1" + local wmma="$2" + local dest="$OUT/$arm" + mkdir -p "$dest/trace" "$dest/art" + env \ + VT_785_P1_GPU=1 \ + VT_785_P1_OUT="$dest/art" \ + VT_ATTN_PREFILL_SHAREDK_WMMA="$wmma" \ + VT_ATTN_PREFILL_FLASH_SHAREDK=1 \ + VT_ATTN_DECODE_OPT=1 \ + VT_ATTN_DECODE_GQA=1 \ + VT_ROCM_ATTN_CPU_REF=0 \ + VT_CPU_REF=0 \ + "$ROCPROF" --kernel-trace --output-format csv \ + -d "$dest/trace" -o "ktrace" \ + -- "$BIN" --test-case="*product seam*" + local rc=$? + if (( rc != 0 )); then + die "arm $arm binary rc=$rc (fail closed; 77 is skip, not a pass; no retry)" + fi + python3 "$PARSER" "$dest/trace" --expect "$arm" \ + | tee "$dest/trace-class.txt" + local hashes="$dest/art/fixture-hashes.txt" + [[ -f "$hashes" ]] || die "arm $arm missing fixture-hashes.txt" + grep -qx "q_bf16=$QHASH" "$hashes" || die "arm $arm Q hash mismatch" + grep -qx "k_bf16=$KHASH" "$hashes" || die "arm $arm K hash mismatch" + grep -qx "v_bf16=$VHASH" "$hashes" || die "arm $arm V hash mismatch" +} + +run_arm A 1 +run_arm B 0 + +python3 - < +#include +#include +#include +#include +#include +#include + +namespace vt_785_p1 { + +constexpr int64_t kT = 64; +constexpr int64_t kHq = 2; +constexpr int64_t kHk = 1; +constexpr int64_t kD = 256; +constexpr int64_t kBlock = 16; +constexpr float kScale = 0.0625f; +constexpr bool kCausal = true; +constexpr int64_t kWindowLeft = 32; +constexpr int64_t kWindowRight = 0; +constexpr uint32_t kQSeed = 78525601u; +constexpr uint32_t kKSeed = 78525602u; +constexpr uint32_t kVSeed = 78525603u; + +constexpr double kAbsFloor = 1.5e-2; +constexpr double kRel = 1.0e-2; +constexpr double kMinCorr = 0.999; + +constexpr const char* kQHash = + "27f164e220edf8d37ddfd1783f0c390968b9d2314ff5d89b7aea3ef50f0eba72"; +constexpr const char* kKHash = + "d803e64df022e1cb2049b06d65fc36da02a79ac09f77831e967237e47d49f7ca"; +constexpr const char* kVHash = + "910548159fd5a5a478573dac6de96f374014cb0b76dac7389071f265dfe1cae7"; + +inline uint16_t F32ToBf16Bits(float f) { + uint32_t x; + std::memcpy(&x, &f, sizeof(x)); + const uint32_t rounding = 0x7fffu + ((x >> 16) & 1u); + return static_cast((x + rounding) >> 16); +} +inline float Bf16BitsToF32(uint16_t b) { + uint32_t x = static_cast(b) << 16; + float f; + std::memcpy(&f, &x, sizeof(f)); + return f; +} + +inline std::vector RandF32(size_t n, uint32_t seed) { + std::vector v(n); + uint32_t s = seed; + for (auto& x : v) { + s = s * 1664525u + 1013904223u; + x = (static_cast(s >> 8) / static_cast(1u << 24)) * 4.0f - 2.0f; + } + return v; +} + +struct Sha256 { + uint32_t s[8] = {0x6a09e667u, 0xbb67ae85u, 0x3c6ef372u, 0xa54ff53au, + 0x510e527fu, 0x9b05688cu, 0x1f83d9abu, 0x5be0cd19u}; + uint64_t nbits = 0; + uint8_t buf[64]{}; + size_t fill = 0; + + static uint32_t Rotr(uint32_t x, int n) { return (x >> n) | (x << (32 - n)); } + void Block(const uint8_t* p) { + static const uint32_t K[64] = { + 0x428a2f98u, 0x71374491u, 0xb5c0fbcfu, 0xe9b5dba5u, 0x3956c25bu, 0x59f111f1u, + 0x923f82a4u, 0xab1c5ed5u, 0xd807aa98u, 0x12835b01u, 0x243185beu, 0x550c7dc3u, + 0x72be5d74u, 0x80deb1feu, 0x9bdc06a7u, 0xc19bf174u, 0xe49b69c1u, 0xefbe4786u, + 0x0fc19dc6u, 0x240ca1ccu, 0x2de92c6fu, 0x4a7484aau, 0x5cb0a9dcu, 0x76f988dau, + 0x983e5152u, 0xa831c66du, 0xb00327c8u, 0xbf597fc7u, 0xc6e00bf3u, 0xd5a79147u, + 0x06ca6351u, 0x14292967u, 0x27b70a85u, 0x2e1b2138u, 0x4d2c6dfcu, 0x53380d13u, + 0x650a7354u, 0x766a0abbu, 0x81c2c92eu, 0x92722c85u, 0xa2bfe8a1u, 0xa81a664bu, + 0xc24b8b70u, 0xc76c51a3u, 0xd192e819u, 0xd6990624u, 0xf40e3585u, 0x106aa070u, + 0x19a4c116u, 0x1e376c08u, 0x2748774cu, 0x34b0bcb5u, 0x391c0cb3u, 0x4ed8aa4au, + 0x5b9cca4fu, 0x682e6ff3u, 0x748f82eeu, 0x78a5636fu, 0x84c87814u, 0x8cc70208u, + 0x90befffau, 0xa4506cebu, 0xbef9a3f7u, 0xc67178f2u}; + uint32_t w[64]; + for (int i = 0; i < 16; ++i) { + w[i] = (uint32_t(p[4 * i]) << 24) | (uint32_t(p[4 * i + 1]) << 16) | + (uint32_t(p[4 * i + 2]) << 8) | uint32_t(p[4 * i + 3]); + } + for (int i = 16; i < 64; ++i) { + const uint32_t s0 = Rotr(w[i - 15], 7) ^ Rotr(w[i - 15], 18) ^ (w[i - 15] >> 3); + const uint32_t s1 = Rotr(w[i - 2], 17) ^ Rotr(w[i - 2], 19) ^ (w[i - 2] >> 10); + w[i] = w[i - 16] + s0 + w[i - 7] + s1; + } + uint32_t a = s[0], b = s[1], c = s[2], d = s[3], e = s[4], f = s[5], g = s[6], + h = s[7]; + for (int i = 0; i < 64; ++i) { + const uint32_t S1 = Rotr(e, 6) ^ Rotr(e, 11) ^ Rotr(e, 25); + const uint32_t ch = (e & f) ^ ((~e) & g); + const uint32_t t1 = h + S1 + ch + K[i] + w[i]; + const uint32_t S0 = Rotr(a, 2) ^ Rotr(a, 13) ^ Rotr(a, 22); + const uint32_t maj = (a & b) ^ (a & c) ^ (b & c); + const uint32_t t2 = S0 + maj; + h = g; + g = f; + f = e; + e = d + t1; + d = c; + c = b; + b = a; + a = t1 + t2; + } + s[0] += a; + s[1] += b; + s[2] += c; + s[3] += d; + s[4] += e; + s[5] += f; + s[6] += g; + s[7] += h; + } + void Update(const void* data, size_t n) { + const auto* p = static_cast(data); + nbits += uint64_t(n) * 8u; + while (n) { + const size_t take = std::min(n, size_t(64 - fill)); + std::memcpy(buf + fill, p, take); + fill += take; + p += take; + n -= take; + if (fill == 64) { + Block(buf); + fill = 0; + } + } + } + std::string Hex() { + uint8_t tail[64 + 8]; + std::memset(tail, 0, sizeof(tail)); + std::memcpy(tail, buf, fill); + tail[fill] = 0x80; + size_t used = fill + 1; + if (used > 56) { + Block(tail); + std::memset(tail, 0, 64); + } + for (int i = 0; i < 8; ++i) tail[63 - i] = static_cast(nbits >> (8 * i)); + Block(tail); + std::string out(64, '0'); + static const char* hex = "0123456789abcdef"; + for (int i = 0; i < 8; ++i) { + for (int b = 0; b < 4; ++b) { + const uint8_t v = static_cast(s[i] >> (24 - 8 * b)); + out[static_cast(i * 8 + b * 2)] = hex[v >> 4]; + out[static_cast(i * 8 + b * 2 + 1)] = hex[v & 0xf]; + } + } + return out; + } +}; + +inline std::string Sha256U16Le(const std::vector& v) { + Sha256 h; + std::vector raw(v.size() * 2); + for (size_t i = 0; i < v.size(); ++i) { + raw[2 * i] = static_cast(v[i] & 0xff); + raw[2 * i + 1] = static_cast((v[i] >> 8) & 0xff); + } + h.Update(raw.data(), raw.size()); + return h.Hex(); +} + +struct Fixture { + std::vector q_bf16, k_bf16, v_bf16; + std::vector q_f32, k_f32, v_f32; + std::vector block_table, seq_lens, qsl; + int64_t num_blocks = 0; +}; + +inline Fixture MakeFixture() { + Fixture f; + f.num_blocks = (kT + kBlock - 1) / kBlock; + auto q = RandF32(static_cast(kT * kHq * kD), kQSeed); + auto k = RandF32(static_cast(f.num_blocks * kBlock * kHk * kD), kKSeed); + auto v = RandF32(static_cast(f.num_blocks * kBlock * kHk * kD), kVSeed); + f.q_bf16.resize(q.size()); + f.k_bf16.resize(k.size()); + f.v_bf16.resize(v.size()); + f.q_f32.resize(q.size()); + f.k_f32.resize(k.size()); + f.v_f32.resize(v.size()); + for (size_t i = 0; i < q.size(); ++i) { + f.q_bf16[i] = F32ToBf16Bits(q[i]); + f.q_f32[i] = Bf16BitsToF32(f.q_bf16[i]); + } + for (size_t i = 0; i < k.size(); ++i) { + f.k_bf16[i] = F32ToBf16Bits(k[i]); + f.k_f32[i] = Bf16BitsToF32(f.k_bf16[i]); + f.v_bf16[i] = F32ToBf16Bits(v[i]); + f.v_f32[i] = Bf16BitsToF32(f.v_bf16[i]); + } + f.block_table.resize(static_cast(f.num_blocks)); + for (int64_t i = 0; i < f.num_blocks; ++i) + f.block_table[static_cast(i)] = static_cast(i); + f.seq_lens = {static_cast(kT)}; + f.qsl = {0, static_cast(kT)}; + return f; +} + +inline std::vector Oracle(const Fixture& f) { + const int64_t qpk = kHq / kHk; + std::vector out(static_cast(kT * kHq * kD), 0.0f); + for (int64_t local = 0; local < kT; ++local) { + const int64_t p = local; + const int64_t jmin = std::max(0, p - kWindowLeft); + int64_t jmax = kCausal ? p : kT - 1; + jmax = std::min(jmax, p + kWindowRight); + jmax = std::min(jmax, kT - 1); + for (int64_t h = 0; h < kHq; ++h) { + const int64_t g = h / qpk; + const int64_t qoff = (local * kHq + h) * kD; + std::vector sc(static_cast(jmax - jmin + 1)); + float m = -std::numeric_limits::infinity(); + for (int64_t j = jmin; j <= jmax; ++j) { + const int64_t blk = f.block_table[static_cast(j / kBlock)]; + const int64_t off = j % kBlock; + const int64_t kbase = ((blk * kBlock + off) * kHk + g) * kD; + float dot = 0.0f; + for (int64_t e = 0; e < kD; ++e) + dot += f.q_f32[static_cast(qoff + e)] * f.k_f32[static_cast(kbase + e)]; + dot *= kScale; + sc[static_cast(j - jmin)] = dot; + if (dot > m) m = dot; + } + float denom = 0.0f; + for (float& s : sc) { + s = std::exp(s - m); + denom += s; + } + const float inv = 1.0f / denom; + for (int64_t e = 0; e < kD; ++e) { + float a = 0.0f; + for (int64_t j = jmin; j <= jmax; ++j) { + const int64_t blk = f.block_table[static_cast(j / kBlock)]; + const int64_t off = j % kBlock; + const int64_t vbase = ((blk * kBlock + off) * kHk + g) * kD; + a += sc[static_cast(j - jmin)] * inv * f.v_f32[static_cast(vbase + e)]; + } + out[static_cast(qoff + e)] = a; + } + } + } + return out; +} + +struct OracleStats { + double max_abs = 0; + double corr = 0; + size_t nonfinite = 0; + size_t violations = 0; + bool ok = false; +}; + +inline OracleStats Score(const std::vector& got, const std::vector& ref) { + OracleStats s; + if (got.size() != ref.size() || got.empty()) return s; + double mean_g = 0, mean_r = 0; + for (size_t i = 0; i < ref.size(); ++i) { + if (!std::isfinite(got[i]) || !std::isfinite(ref[i])) ++s.nonfinite; + mean_g += got[i]; + mean_r += ref[i]; + } + mean_g /= static_cast(ref.size()); + mean_r /= static_cast(ref.size()); + double num = 0, dg = 0, dr = 0; + for (size_t i = 0; i < ref.size(); ++i) { + const double err = std::abs(got[i] - ref[i]); + if (err > s.max_abs) s.max_abs = err; + const double thr = kAbsFloor + kRel * std::abs(static_cast(ref[i])); + if (err > thr) ++s.violations; + const double ag = got[i] - mean_g; + const double ar = ref[i] - mean_r; + num += ag * ar; + dg += ag * ag; + dr += ar * ar; + } + s.corr = (dg > 0.0 && dr > 0.0) ? (num / std::sqrt(dg * dr)) : 0.0; + s.ok = s.nonfinite == 0 && s.violations == 0 && s.corr >= kMinCorr; + return s; +} + +inline constexpr const char* kExactWmma[] = { + "PagedAttnPrefillSharedKWmma<2, 8, 16, 32, false>", + "PagedAttnPrefillSharedKWmma<2,8,16,32,false>", + "PagedAttnPrefillSharedKWmmaILi2ELi8ELi16ELi32ELb0E", +}; +inline constexpr const char* kExactScalar[] = { + "PagedAttnPrefillSharedK<2, 8, 32, 32>", + "PagedAttnPrefillSharedK<2,8,32,32>", + "PagedAttnPrefillSharedKILi2ELi8ELi32ELi32EE", +}; + +inline std::string ScalarLines(const std::string& text) { + std::string filtered; + filtered.reserve(text.size()); + size_t i = 0; + while (i < text.size()) { + const size_t nl = text.find('\n', i); + const size_t end = nl == std::string::npos ? text.size() : nl; + const auto line = text.substr(i, end - i); + if (line.find("SharedKWmma") == std::string::npos) { + filtered.append(line); + filtered.push_back('\n'); + } + i = end == text.size() ? end : end + 1; + } + return filtered; +} + +inline bool HasAny(const std::string& text, const char* const* ms, size_t n) { + for (size_t i = 0; i < n; ++i) { + if (text.find(ms[i]) != std::string::npos) return true; + } + return false; +} + +inline std::string StripExact(const std::string& text, const char* const* ms, size_t n) { + std::string out = text; + for (size_t i = 0; i < n; ++i) { + for (;;) { + const auto pos = out.find(ms[i]); + if (pos == std::string::npos) break; + out.erase(pos, std::strlen(ms[i])); + } + } + return out; +} + +// Family vs 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 => '?'. +inline char ClassifyArm(const std::string& text) { + const auto scalar_blob = ScalarLines(text); + const bool exact_wmma = + HasAny(text, kExactWmma, sizeof(kExactWmma) / sizeof(kExactWmma[0])); + const bool exact_scalar = HasAny( + scalar_blob, kExactScalar, sizeof(kExactScalar) / sizeof(kExactScalar[0])); + const bool wmma_family = text.find("PagedAttnPrefillSharedKWmma") != std::string::npos; + const bool scalar_family = + scalar_blob.find("PagedAttnPrefillSharedK") != std::string::npos; + const bool other_wmma = + StripExact(text, kExactWmma, sizeof(kExactWmma) / sizeof(kExactWmma[0])) + .find("PagedAttnPrefillSharedKWmma") != std::string::npos; + const bool other_scalar = + StripExact(scalar_blob, kExactScalar, + sizeof(kExactScalar) / sizeof(kExactScalar[0])) + .find("PagedAttnPrefillSharedK") != std::string::npos; + if (exact_wmma && !scalar_family && !other_wmma) return 'A'; + if (exact_scalar && !wmma_family && !other_scalar) return 'B'; + return '?'; +} + +} // namespace vt_785_p1 diff --git a/tests/vt/test_ops_paged_attn_sharedk_wmma_p1.cpp b/tests/vt/test_ops_paged_attn_sharedk_wmma_p1.cpp new file mode 100644 index 000000000..210a28ba7 --- /dev/null +++ b/tests/vt/test_ops_paged_attn_sharedk_wmma_p1.cpp @@ -0,0 +1,122 @@ +// #785 P1 host package: frozen fixture hashes, oracle invariants, +// exclusive trace classifier, CTest-registration lock. CPU-only. +#include + +#include +#include +#include +#include +#include +#include + +#include "vt/sharedk_wmma_p1_fixture.h" + +using namespace vt_785_p1; + +TEST_CASE("P1 fixture freeze: shapes, seeds, hashes") { + CHECK(kT >= 64); + CHECK(kHq / kHk == 2); + CHECK(kD == 256); + const auto f = MakeFixture(); + CHECK(f.q_bf16.size() == static_cast(kT * kHq * kD)); + CHECK(f.k_bf16.size() == static_cast(f.num_blocks * kBlock * kHk * kD)); + CHECK(f.seq_lens.size() == 1); + CHECK(f.seq_lens[0] == kT); + CHECK(f.qsl.back() == kT); + CHECK(Sha256U16Le(f.q_bf16) == kQHash); + CHECK(Sha256U16Le(f.k_bf16) == kKHash); + CHECK(Sha256U16Le(f.v_bf16) == kVHash); +} + +TEST_CASE("P1 host f32 oracle freeze") { + const auto f = MakeFixture(); + const auto ref = Oracle(f); + CHECK(ref.size() == f.q_f32.size()); + double max_abs = 0, mean = 0; + for (float x : ref) { + CHECK(std::isfinite(x)); + max_abs = std::max(max_abs, static_cast(std::abs(x))); + mean += x; + } + mean /= static_cast(ref.size()); + CHECK(max_abs > 1.0); + CHECK(max_abs < 4.0); + CHECK(std::abs(mean) < 0.05); +} + +TEST_CASE("P1 BF16 tolerance: rounded-ref is GREEN; swapped is RED") { + const auto f = MakeFixture(); + const auto ref = Oracle(f); + CHECK(Score(ref, ref).ok); + auto swapped = ref; + if (!swapped.empty()) swapped[0] = -swapped[0] - 10.0f; + CHECK_FALSE(Score(swapped, ref).ok); +} + +TEST_CASE("P1 zero-variance / empty fail closed") { + CHECK_FALSE(Score({}, {1.0f}).ok); + CHECK_FALSE(Score({1.0f}, {}).ok); + const std::vector z(8, 0.0f); + auto s = Score(z, z); + CHECK(s.nonfinite == 0); + CHECK(s.corr == doctest::Approx(0.0)); + CHECK_FALSE(s.ok); +} + +TEST_CASE("P1 exclusive classifier: exact spec XOR; wrong family UNKNOWN") { + const std::string a = + "Dispatch,Kernel_Name\n" + "0,\"void vt::rocm::(anonymous namespace)::" + "PagedAttnPrefillSharedKWmma<2, 8, 16, 32, false>(__hip_bfloat16*)\"\n"; + const std::string b = + "Dispatch,Kernel_Name\n" + "0,\"void vt::rocm::(anonymous namespace)::" + "PagedAttnPrefillSharedK<2, 8, 32, 32>(__hip_bfloat16*)\"\n"; + const std::string mangled_a = + "_ZN2vt4rocm12_GLOBAL__N_127PagedAttnPrefillSharedKWmmaILi2ELi8ELi16ELi32ELb0EEEv"; + const std::string mangled_b = + "_ZN2vt4rocm12_GLOBAL__N_124PagedAttnPrefillSharedKILi2ELi8ELi32ELi32EEEv"; + const std::string empty = "Dispatch,Kernel_Name\n"; + const std::string wrong_scalar = + "PagedAttnPrefillSharedK<2, 8, 16, 32>"; + const std::string d512_scalar = + "PagedAttnPrefillSharedK<2, 16, 16, 16>"; + const std::string wrong_wmma = + "PagedAttnPrefillSharedKWmma<2, 8, 16, 16, false>"; + CHECK(ClassifyArm(a) == 'A'); + CHECK(ClassifyArm(b) == 'B'); + CHECK(ClassifyArm(mangled_a) == 'A'); + CHECK(ClassifyArm(mangled_b) == 'B'); + CHECK(ClassifyArm(empty) == '?'); + CHECK(ClassifyArm(a + b) == '?'); + CHECK(ClassifyArm(wrong_scalar) == '?'); + CHECK(ClassifyArm(d512_scalar) == '?'); + CHECK(ClassifyArm(wrong_wmma) == '?'); + CHECK(ClassifyArm(std::string("PagedAttnPrefillSharedKILi2ELi8ELi16ELi32EE")) == '?'); +} + +TEST_CASE("P1 eligibility pins SharedK d=256 qg=2 T>=64 one request") { + CHECK(kT >= 64); + CHECK(kD == 256); + CHECK(kHq == 2); + CHECK(kHk == 1); + CHECK(kCausal); + CHECK(kScale == doctest::Approx(1.0f / std::sqrt(256.0f))); + CHECK(kWindowLeft == 32); + CHECK(kWindowRight == 0); +} + +TEST_CASE("P1 GPU binary is not registered as ordinary CTest") { + std::string path = __FILE__; + const auto slash = path.find_last_of('/'); + REQUIRE(slash != std::string::npos); + const std::string cmake = path.substr(0, slash) + "/../CMakeLists.txt"; + std::ifstream in(cmake); + REQUIRE(in); + std::string text((std::istreambuf_iterator(in)), std::istreambuf_iterator()); + CHECK(text.find("test_ops_paged_attn_sharedk_wmma_p1_gpu") != std::string::npos); + CHECK(text.find("add_test(NAME test_ops_paged_attn_sharedk_wmma_p1_gpu") == + std::string::npos); + CHECK(text.find("vllm_cpp_add_test(test_ops_paged_attn_sharedk_wmma_p1_gpu") == + std::string::npos); +} diff --git a/tests/vt/test_ops_paged_attn_sharedk_wmma_p1_gpu.cpp b/tests/vt/test_ops_paged_attn_sharedk_wmma_p1_gpu.cpp new file mode 100644 index 000000000..4ac2d478a --- /dev/null +++ b/tests/vt/test_ops_paged_attn_sharedk_wmma_p1_gpu.cpp @@ -0,0 +1,186 @@ +// #785 P1 GPU product-seam witness. Calls vt::PagedAttention on kROCM. +// Does NOT launch PagedAttnPrefillSharedKWmma directly. +// Not registered with CTest. Run only via tests/scripts/run-785-p1.sh +// after Researcher GPU GO. +// +// Missing VT_785_P1_GPU, missing device, env conflict, or hash mismatch +// => exit 5 (fail closed). Runner treats 77/nonzero as P1 failure. +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vt/sharedk_wmma_p1_fixture.h" +#include "vt/backend.h" +#include "vt/dtype.h" +#include "vt/ops.h" +#include "vt/rocm/rocm_runtime.h" + +using vt::Backend; +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::PagedAttentionArgs; +using vt::Queue; +using vt::Tensor; +using namespace vt_785_p1; + +namespace { + +[[noreturn]] void FailClosed(const char* why) { + std::fprintf(stderr, "\n*** P1 GPU FAIL CLOSED (exit 5) ***\n%s\n", why); + std::exit(5); +} + +const char* EnvOr(const char* key) { + const char* e = std::getenv(key); + return e ? e : ""; +} + +bool EnvIs1(const char* key) { + const char* e = std::getenv(key); + return e != nullptr && e[0] == '1' && e[1] == '\0'; +} +bool EnvIs0(const char* key) { + const char* e = std::getenv(key); + return e != nullptr && e[0] == '0' && e[1] == '\0'; +} + +Tensor Contig(void* data, DType dt, Device dev, const std::vector& shape) { + Tensor t; + t.data = data; + t.dtype = dt; + t.device = dev; + t.rank = static_cast(shape.size()); + int64_t stride = 1; + for (int i = t.rank - 1; i >= 0; --i) { + t.shape[i] = shape[static_cast(i)]; + t.stride[i] = stride; + stride *= shape[static_cast(i)]; + } + return t; +} + +struct DeviceBuf { + Backend& b; + void* p = nullptr; + size_t bytes = 0; + Tensor t; + DeviceBuf(Backend& b_, Queue& q, DType dt, const std::vector& shape, const void* host) + : b(b_) { + int64_t n = 1; + for (auto s : shape) n *= s; + bytes = static_cast(n) * vt::SizeOf(dt); + p = b.Alloc(bytes == 0 ? 1 : bytes); + if (host != nullptr) b.Copy(q, p, host, bytes); + t = Contig(p, dt, Device{DeviceType::kROCM, 0}, shape); + } + ~DeviceBuf() { + if (p) b.Free(p); + } + void Download(Queue& q, void* dst) { + b.Copy(q, dst, p, bytes); + b.Synchronize(q); + } +}; + +} // namespace + +TEST_CASE("P1 product seam vt::PagedAttention d=256 SharedK") { + if (std::getenv("VT_785_P1_GPU") == nullptr) { + FailClosed("VT_785_P1_GPU unset"); + } + if (!vt::rocm::DeviceAvailable()) { + FailClosed("no ROCm device"); + } + const char* outdir = std::getenv("VT_785_P1_OUT"); + if (outdir == nullptr || outdir[0] == '\0') { + FailClosed("VT_785_P1_OUT unset"); + } + + if (EnvIs1("VT_ROCM_ATTN_CPU_REF") || EnvIs1("VT_CPU_REF")) { + FailClosed("CPU-ref must be off"); + } + if (EnvIs0("VT_ATTN_DECODE_OPT") || EnvIs0("VT_ATTN_DECODE_GQA") || + EnvIs0("VT_ATTN_PREFILL_FLASH_SHAREDK")) { + FailClosed("decode-opt / decode-GQA / SharedK must be on"); + } + if (!EnvIs1("VT_ATTN_DECODE_OPT") || !EnvIs1("VT_ATTN_DECODE_GQA") || + !EnvIs1("VT_ATTN_PREFILL_FLASH_SHAREDK")) { + FailClosed("decode-opt / decode-GQA / SharedK must be explicitly 1"); + } + const bool wmma1 = EnvIs1("VT_ATTN_PREFILL_SHAREDK_WMMA"); + const bool wmma0 = EnvIs0("VT_ATTN_PREFILL_SHAREDK_WMMA"); + if (wmma1 == wmma0) { + FailClosed("VT_ATTN_PREFILL_SHAREDK_WMMA must be explicit 1 or 0"); + } + + const auto f = MakeFixture(); + const std::string qh = Sha256U16Le(f.q_bf16); + const std::string kh = Sha256U16Le(f.k_bf16); + const std::string vh = Sha256U16Le(f.v_bf16); + if (qh != kQHash || kh != kKHash || vh != kVHash) { + FailClosed("fixture Q/K/V SHA-256 mismatch vs preregistered"); + } + { + std::ofstream h(std::string(outdir) + "/fixture-hashes.txt"); + if (!h) FailClosed("cannot write fixture-hashes.txt"); + h << "q_bf16=" << qh << "\nk_bf16=" << kh << "\nv_bf16=" << vh << "\n"; + } + { + std::ofstream e(std::string(outdir) + "/env.txt"); + if (!e) FailClosed("cannot write env.txt"); + e << "VT_785_P1_GPU=" << EnvOr("VT_785_P1_GPU") << "\n" + << "VT_ATTN_PREFILL_SHAREDK_WMMA=" << EnvOr("VT_ATTN_PREFILL_SHAREDK_WMMA") << "\n" + << "VT_ATTN_PREFILL_FLASH_SHAREDK=" << EnvOr("VT_ATTN_PREFILL_FLASH_SHAREDK") << "\n" + << "VT_ATTN_DECODE_OPT=" << EnvOr("VT_ATTN_DECODE_OPT") << "\n" + << "VT_ATTN_DECODE_GQA=" << EnvOr("VT_ATTN_DECODE_GQA") << "\n" + << "VT_ROCM_ATTN_CPU_REF=" << EnvOr("VT_ROCM_ATTN_CPU_REF") << "\n" + << "VT_CPU_REF=" << EnvOr("VT_CPU_REF") << "\n"; + } + + const auto ref = Oracle(f); + Backend& rocm = vt::GetBackend(DeviceType::kROCM); + Queue q = rocm.CreateQueue(); + DeviceBuf dq(rocm, q, DType::kBF16, {kT, kHq, kD}, f.q_bf16.data()); + DeviceBuf dk(rocm, q, DType::kBF16, {f.num_blocks, kBlock, kHk, kD}, f.k_bf16.data()); + DeviceBuf dv(rocm, q, DType::kBF16, {f.num_blocks, kBlock, kHk, kD}, f.v_bf16.data()); + DeviceBuf dbt(rocm, q, DType::kI32, {1, f.num_blocks}, f.block_table.data()); + DeviceBuf dsl(rocm, q, DType::kI32, {1}, f.seq_lens.data()); + DeviceBuf dqsl(rocm, q, DType::kI32, {2}, f.qsl.data()); + DeviceBuf dout(rocm, q, DType::kBF16, {kT, kHq, kD}, nullptr); + + PagedAttentionArgs args{kScale, /*causal=*/true}; + args.window_size = vt::AttentionWindow{static_cast(kWindowLeft), + static_cast(kWindowRight)}; + args.query_start_loc_host = f.qsl.data(); + args.max_seq_len = static_cast(kT); + vt::PagedAttention(q, dout.t, dq.t, dk.t, dv.t, dbt.t, dsl.t, dqsl.t, args); + + std::vector got(f.q_bf16.size(), 0); + dout.Download(q, got.data()); + std::vector got_f(got.size()); + for (size_t i = 0; i < got.size(); ++i) got_f[i] = Bf16BitsToF32(got[i]); + const auto st = Score(got_f, ref); + + { + std::ofstream o(std::string(outdir) + "/out.bf16", std::ios::binary); + if (!o) FailClosed("cannot write out.bf16"); + o.write(reinterpret_cast(got.data()), + static_cast(got.size() * sizeof(uint16_t))); + } + { + std::ofstream m(std::string(outdir) + "/metrics.txt"); + if (!m) FailClosed("cannot write metrics.txt"); + m << "max_abs=" << st.max_abs << "\ncorr=" << st.corr << "\nnonfinite=" << st.nonfinite + << "\nviolations=" << st.violations << "\noracle_ok=" << (st.ok ? 1 : 0) + << "\nn=" << got.size() << "\n"; + } + if (!st.ok) FailClosed("oracle bar miss"); +} diff --git a/tests/vt/test_rocm_arch.cpp b/tests/vt/test_rocm_arch.cpp index 6cd8b7363..8531164f3 100644 --- a/tests/vt/test_rocm_arch.cpp +++ b/tests/vt/test_rocm_arch.cpp @@ -72,6 +72,26 @@ TEST_CASE("major outside [9, 12] is declined") { CHECK_FALSE(CapabilityFromGcnArch("gfx1300").has_value()); // no such generation } +TEST_CASE("SharedK WMMA host gate is gfx1200/gfx1201 prefix, not substring") { + using vt::rocm::GcnArchNameIsGfx12PrefillWmma; + CHECK(GcnArchNameIsGfx12PrefillWmma("gfx1200")); + CHECK(GcnArchNameIsGfx12PrefillWmma("gfx1201")); + CHECK(GcnArchNameIsGfx12PrefillWmma("gfx1201:xnack-")); + CHECK(GcnArchNameIsGfx12PrefillWmma("gfx1201:sramecc+")); + CHECK(GcnArchNameIsGfx12PrefillWmma("gfx1200:xnack-")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("gfx1100")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("gfx1202")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("gfx1210")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("foogfx1201")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("agfx1201")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("gfx12010")); + CHECK_FALSE(GcnArchNameIsGfx12PrefillWmma("gfx120")); + static_assert(GcnArchNameIsGfx12PrefillWmma("gfx1201:xnack-")); + static_assert(!GcnArchNameIsGfx12PrefillWmma("foogfx1201")); + static_assert(!GcnArchNameIsGfx12PrefillWmma("gfx12010")); +} + TEST_CASE("the parse is constexpr, so a wrong answer is a compile error") { // Not decoration: it is what lets the capability be asserted without a device. static_assert(CapabilityFromGcnArch("gfx1100")->first == 11);