From 2ba031bf3bc34e849d0dd3d42df3fb0bbbdd501b Mon Sep 17 00:00:00 2001 From: Don Mirror Date: Fri, 14 Aug 2026 17:56:08 -0500 Subject: [PATCH 1/7] docs(rocm/gemma4): spec #837 GetBlas dual-slot (split repair) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Independent row/ROCM-GEMMA4-GETBLAS-DUALSLOT from origin/main 3ce5a1dc. Pins donor GetBlas bytes (SHA256 9df2b163…) — dirty 2bb4bd8a is not a clean donor. Hypothesis B only; cause of T=2029 wedge unconfirmed. Host lifetime seam + 0→1→0 / 1→0→1 mutations required (64cb #5). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes] --- .../CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md | 5 + .../evidence/rocm-gemma4-getblas/MANIFEST.md | 18 +++ .../rocm-gemma4-getblas/getblas-fn-67-99.txt | 33 +++++ .agents/issue-index.md | 1 + .agents/specs/rocm-gemma4-getblas-dualslot.md | 122 ++++++++++++++++++ 5 files changed, 179 insertions(+) create mode 100644 .agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md create mode 100644 .agents/evidence/rocm-gemma4-getblas/MANIFEST.md create mode 100644 .agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt create mode 100644 .agents/specs/rocm-gemma4-getblas-dualslot.md diff --git a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md new file mode 100644 index 000000000..859950519 --- /dev/null +++ b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md @@ -0,0 +1,5 @@ +# CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT + +| Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `SPIKE` spec-only | 2026-08-14 — `64cb` split repair; donor slice hashed | diff --git a/.agents/evidence/rocm-gemma4-getblas/MANIFEST.md b/.agents/evidence/rocm-gemma4-getblas/MANIFEST.md new file mode 100644 index 000000000..dc69a77a3 --- /dev/null +++ b/.agents/evidence/rocm-gemma4-getblas/MANIFEST.md @@ -0,0 +1,18 @@ +# Donor evidence — #837 GetBlas dual-slot + +Pinned **bytes**, not a dirty-tree HEAD. Implementation must copy this slice (or a later +immutable replacement that research re-reviews), not re-read `/home/don/llms/vllm.cpp`. + +| Field | Value | +|---|---| +| Donor tree | `/home/don/llms/vllm.cpp` | +| Donor git HEAD | `2bb4bd8a` (dirty; this slice is **uncommitted** on that tree) | +| File | `src/vt/rocm/rocm_matmul_hipblaslt.hip` | +| Lines | 67–99 (`GetBlas`) | +| Slice | `getblas-fn-67-99.txt` | +| SHA256 | `9df2b163bc817db0d9545570136666c8e07a0bb600a01e50288a8f78c4148c51` | +| Recipient | `origin/main` `3ce5a1dc` `src/vt/rocm/rocm_matmul_hipblaslt.hip:72-99` (single `static thread_local Tls tls`) | +| Captured | 2026-08-14 | + +`sha256sum` of the slice file must match the table. Do not treat `2bb4bd8a` as a clean +donor commit. diff --git a/.agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt b/.agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt new file mode 100644 index 000000000..6abf5cc0d --- /dev/null +++ b/.agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt @@ -0,0 +1,33 @@ +// engine worker. Avoid global mutex + SetStream on every GEMM (100s×/token). +BlasCtx GetBlas(int device, hipStream_t stream) { + struct Tls { + int dev = -1; + hipStream_t stream = nullptr; + hipblasHandle_t handle = nullptr; + }; + // Per-device: layer-split pipe must not hipblasDestroy GPU0 while GPU0 GEMMs + // are still queued (same class as SameTls / LayerTls). + static thread_local Tls tls_slots[2]; + Tls& tls = tls_slots[(device == 1) ? 1 : 0]; + // Bind current device outside capture so create/setStream see the right GPU + // after peer-MoE left current on the expert device. + if (!StreamIsCapturing(stream)) { + int cur = -1; + if (hipGetDevice(&cur) != hipSuccess || cur != device) (void)hipSetDevice(device); + } + if (tls.handle == nullptr || tls.dev != device) { + if (tls.handle) { + (void)hipblasDestroy(tls.handle); + tls.handle = nullptr; + } + if (!StreamIsCapturing(stream)) (void)hipSetDevice(device); + CheckBlas(hipblasCreate(&tls.handle), "hipblasCreate"); + tls.dev = device; + tls.stream = nullptr; + } + if (tls.stream != stream) { + CheckBlas(hipblasSetStream(tls.handle, stream), "hipblasSetStream"); + tls.stream = stream; + } + return BlasCtx{tls.handle}; +} diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 0f7074c35..9054038b6 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -325,3 +325,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1068](https://github.com/mudler/vllm.cpp/issues/1068) | `ENG-RELEASE-WINDOWS` | `main` stopped compiling under MSVC at `e34d71379` (#1054), which dropped the `[&kRequired]` capture from the `refuse` lambda in `qwen3_5_weights.cpp` as "the redundant namespace-scope capture". `kMoeExpertLayoutHelp` (`:894`) is namespace-scope and needs no capture; `kRequired` (`:929`) is a function-local `const std::string&` bound to it and IS odr-used in the lambda body, so MSVC rejects it (`error C3493`). Fixed by naming the namespace-scope constant inside the lambda, which satisfies MSVC and keeps the AppleClang diagnostic #1054 removed. It landed green because the guarding gate is a source-TEXT assertion ("rejects `[&kRequired]`, finds `[]`") that passes whether or not the TU compiles, and because `windows-msvc-*` are skipped on `main` (#503) so no baseline existed to regress. An instance of [#503](https://github.com/mudler/vllm.cpp/issues/503) | bug | | [#1092](https://github.com/mudler/vllm.cpp/issues/1092) | `LTX25-GUIDED-VIDEO` | The LTX-2.5 **video** denoise loop runs one UNGUIDED forward per step: `Ltx2PhaseRecipe::video_guidance` is set by every recipe (`src/vllm/model_executor/models/ltx2_pipeline.cpp:1069 @ b5756ea8c`) and read by nothing, so a `pipeline_kind = one_stage` render ignores `cfg_scale = 3.0`, `stg_scale = 1.0`, `rescale_scale = 0.7` and `modality_scale = 3.0` and denoises along a different trajectory than `ti2vid_one_stage.py:221-226 @ fd4ded7f`, which builds a `FactoryGuidedDenoiser` from exactly those. `allow_guidance_override` (`ltx2_pipeline.h:534`) is dead the same way. Positive control for the grep: the same command for `audio_guidance` returns the T2A consumer at `ltx2_video.cpp:3527`. Blocks four more pipelines on one missing seam (`a2vid_two_stage.py:230`, `ti2vid_two_stages.py:248`, `ti2vid_two_stages_hq.py:271`, `keyframe_interpolation.py:232`). Spec [`ltx25-guided-video.md`](specs/ltx25-guided-video.md) | bug | | [#1111](https://github.com/mudler/vllm.cpp/issues/1111) | `LTX25-GUIDED-VIDEO` | `Ltx2T2aGuidedStep` (`src/vllm/model_executor/models/ltx2_t2a.cpp:203-214` @ `420f6b474`) builds the STG block mask and refuses when NO bit is set, and that one predicate catches two different things. One is a request that disagrees with the CHECKPOINT - `[28]` on a two-block DiT - which is a LOCAL condition, because upstream only ever runs 48-block checkpoints and this port runs reduced ones. The other is an EMPTY list, which is upstream's OWN documented spelling for "STG off" and is refused here. Measured at Lightricks/LTX-2 `fd4ded7f`: `ltx-pipelines/docs/multimodal-guidance.md:13` says "Set to `[]` to disable STG", in the same table and the same idiom as `stg_scale` -> 0.0 and `cfg_scale` -> 1.0; `MultiModalGuiderParams.stg_blocks` DEFAULTS to `[]` (`guiders.py:204`, `field(default_factory=list)`); `--audio-stg-blocks` is `nargs="*"` (`args.py:1039-1045`, `:1107-1113`) so the empty list has a CLI spelling, and `nargs="+"` was the one-character way to forbid it; `LTX_2_3_HQ_PARAMS` SHIPS `stg_blocks=[]` on both modalities (`constants.py:105,:113`); and a whole-tree search found NO validation of `stg_blocks` anywhere - no emptiness, length or range check, no assert, no raise - with the null results recorded. `blocks=None` means EVERY block and `blocks=[]` means NO block (`perturbations.py:26-33`), so the empty list is how a caller says the second thing, and both this tree's parsers go out of their way to keep PRESENT-and-empty distinct from ABSENT - then refuse the one configuration that distinction exists to express. Upstream does not skip the pass either: `do_perturbed_generation` reads `stg_scale` alone (`guiders.py:279-281`), so an empty list disables the STG SIGNAL and not the STG COST. FOUND while repairing the fresh review of [#1092](https://github.com/mudler/vllm.cpp/issues/1092), whose B3 finding is the same divergence on the VIDEO path; that half IS fixed in flow (`ltx2_video.cpp` `ApplyGuidanceOverrides`, `ltx2_denoisers.cpp` `check_reaches_a_block`). NOT fixed here because this half landed with [#1032](https://github.com/mudler/vllm.cpp/issues/1032) and moving it changes a landed row's gated behaviour and its case "a perturbed pass over NO block is refused", which is the normal row-spec-review path rather than an in-flow repair. Until then the two arms disagree about the same request. Listed under `## Owed` in [`ltx25-guided-video.md`](specs/ltx25-guided-video.md) | bug | +| [#837](https://github.com/mudler/vllm.cpp/issues/837) | `ROCM-GEMMA4-GETBLAS-DUALSLOT` | ROCm: dual-slot GetBlas TLS + host lifetime so two HIP streams do not share one handle | bug | diff --git a/.agents/specs/rocm-gemma4-getblas-dualslot.md b/.agents/specs/rocm-gemma4-getblas-dualslot.md new file mode 100644 index 000000000..c676493af --- /dev/null +++ b/.agents/specs/rocm-gemma4-getblas-dualslot.md @@ -0,0 +1,122 @@ +# Spec: ROCm GetBlas dual-slot TLS (Gemma-4 peer-MoE) + +- **Issue:** https://github.com/mudler/vllm.cpp/issues/837 +- **Row slug:** `ROCM-GEMMA4-GETBLAS-DUALSLOT` — child of `BACKEND-ROCM` (#41). Not a new KERNEL family. Separate from #697 / `KERNEL-ROCM-GEMMA4-BC64-FA-PREFILL`. +- **Worktree / branch (this unit only):** `/home/don/llms/vllm.cpp-getblas` · `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` +- **Base / recipient:** `origin/main` `3ce5a1dc` `src/vt/rocm/rocm_matmul_hipblaslt.hip:72-99` +- **Donor bytes:** `.agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt` SHA256 `9df2b163bc817db0d9545570136666c8e07a0bb600a01e50288a8f78c4148c51` (dirty lab `/home/don/llms/vllm.cpp` HEAD `2bb4bd8a` **plus uncommitted**; HEAD is not a clean donor). +- **Implementer:** hermes-vllm (lab). **Reviewer:** research (spec then impl). **Operator/smoke:** coord-help-20260812. +- **Git:** spec-only first (coord `25c9` / research `5071` / BLOCK `64cb`); implementation is a later commit **on this same row branch** after spec GREEN. Independent RED/GREEN from #838/#839. One PR per row. No shared `row/ROCM-GEMMA4-XDEV-MOE` landing history. +- **Supersedes for review:** `20332292` (BLOCK) and preview `c4fbe6e9` (not spec-GREEN). Those SHAs lived on a combined branch and are not review targets. + +## Now + +`SPIKE` — spec for review. No product code in this commit. + +**Not a confirmed fix.** `9772` is an accumulation failure class. This row is hypothesis (B) only. Route observed: T≥64 prefill-batch peer path calls `MatmulBT` → `GetBlas` on the expert queue. Cause (single-TLS destroy/create vs cache vs Launch/Finish) is **unconfirmed**. Do not label this a common root with #838 or #839. + +## Upstream / source of the port + +vLLM has no hipBLAS TLS. Source is the pinned donor slice above, cited as a **file:line + hash**, not a tree transplant. + +| Tree | `GetBlas` | Shape | +|---|---|---| +| `origin/main` `3ce5a1dc` | `src/vt/rocm/rocm_matmul_hipblaslt.hip:72-99` | `static thread_local Tls tls`; on `tls.dev != device` → `hipblasDestroy` + `hipblasCreate` | +| hanging `vllm.cpp-bc64fa-r2` `1b1baf43` | `:72-99` | same single TLS | +| donor slice (hash above) | lab `:67-99` | `static thread_local Tls tls_slots[2]`; `Tls& tls = tls_slots[(device == 1) ? 1 : 0]` | + +Callers (`MatmulBT` / `MatmulBTAlphaBeta` / Lt variants in the same file) stay unchanged. T≥64 prefill-batch peer path reaches `GetBlas` via `vt::MatmulBT` on the expert queue. Serial T=19 / `ExpertGeGLUFp8TopKM1` does **not**. + +## Symptom this row owns + +Coord `9772` / `25c9`: 274 matched `moe_prefill_peer_helper` BEGIN/END, then accumulation wedge. Hypothesis (B): single-TLS destroy/create on the hop is shared-resource churn for the **batch** path. This row does not claim to fix T=19 and does not claim T=2029 will generate from this change alone. + +## Scope + +Replace only the TLS storage in `GetBlas`: + +```c +static thread_local Tls tls_slots[2]; +Tls& tls = tls_slots[(device == 1) ? 1 : 0]; +``` + +Keep the existing capture/`hipGetDevice`/`hipSetDevice`/`hipblasSetStream` body **inside the selected slot**. Device ids other than 0/1 share slot 0 (same as donor). + +## Out of scope + +- Launch/Finish / `PeerPipeTls` / `DequantCacheSlotFor` (#839). +- Indexed T<63 routing (#838). +- hipBLASLt product default, FP8×FP8 Lt, `VT_ROCM_HIPBLASLT`. +- `#697` / any edit of `rocm_paged_attn.hip`. +- Diagnostic `STAGE_SYNC` / `PREFILL_TRACE`. +- More than two slots; devices ≥2. + +## Design + +1. Two process-lifetime per-thread handles. Hop 0→1 must not destroy GPU0's handle while GPU0 GEMMs may still be queued (donor comment at `:74-75`). +2. Create still happens lazily per slot on first use. +3. Stream bind remains per-slot (`tls.stream != stream` → `hipblasSetStream`). +4. Default ON inside `GetBlas` only. No new env. No behavior change outside hipBLAS handle lifetime on device 0/1 hops. +5. **Lifetime invariants:** + - slot[i] handle is destroyed only when that slot is recreated for a **different** `tls.dev` than its index (should not happen if index is a function of `device`); + - hopping 0→1 must leave slot[0].handle live (same pointer identity); + - hopping 1→0 must leave slot[1].handle live; + - `hipSetDevice` after a hop restores `device` before create/setStream; + - capture path still skips setDevice. + +## Risks + +- Slot index `device==1` is a two-GPU lab assumption. A third visible device aliases slot 0. Named; do not invent a map. +- Does not by itself make T=2029 generate. Land before or with #839 as a **separate** immutable head. + +## Tests + +Text search for `tls_slots[2]` / `(device == 1)` is **not** sufficient (research `64cb` stop-ship 5). Tests must observe the new lifetime guarantees. + +### Host load-bearing seam (required) + +Extract or wrap: + +1. `GetBlasSlotIndex(int device) -> 0|1` (`device==1 ? 1 : 0`). +2. A test-only lifecycle recorder (fake `hipblasCreate`/`Destroy`/`SetStream`/`SetDevice`, or a friend/hook compiled into `tests/vt/test_ops_getblas_dualslot.cpp`) that records, per slot: create count, destroy count, last handle identity, last bound stream, last `hipSetDevice` argument, whether the call was under capture. + +Host cases (no GPU required if the seam is fakeable): + +| Case | Expect | +|---|---| +| first use dev0 | slot0 create==1, slot1 create==0 | +| then hop 0→1 | slot0 destroy==0, slot1 create==1, slot0 handle identity unchanged | +| then hop 1→0 | slot1 destroy==0, slot0 create still 1 (no recreate), slot1 handle identity unchanged | +| then hop 1→0→1 | both handles survive; no extra destroy | +| stream change on slot0 | `SetStream` on slot0 only; slot1 stream untouched | +| capture path | no `hipSetDevice` | + +RED mutations (must fail the table): + +- swapped selector (`device==0` → slot 1); +- destroy-on-hop (old single-TLS `if (tls.dev != device) hipblasDestroy`); +- missing stream rebind (`tls.stream != stream` branch deleted); +- capture-path `setDevice` (setDevice runs even when `StreamIsCapturing`). + +### Coord GPU probe (mandatory on impl, not this spec commit) + +On dual visible devices: call `GetBlas(0,s0)`, `GetBlas(1,s1)`, `GetBlas(0,s0)`, `GetBlas(1,s1)` and assert handle pointer identities: 0→1→0 keeps the first GPU0 handle; 1→0→1 keeps the first GPU1 handle. Skip only when `HIP_VISIBLE_DEVICES` empty. This does **not** replace the host seam. + +## Gates + +- Host seam table GREEN without a GPU. +- T=1 decode + Paris + arith on the post-impl binary unchanged vs pre-change KEEP class. +- Operator A/B (`5071`): this is **B**. Run independently (or after A) on **T=2029**; do not bundle with C. p42k only after the smallest passing set. +- `#697` files untouched (`git diff` must not list `rocm_paged_attn.hip` or bc64 tests). +- Default path outside Gemma-4 FP8 xdev `GetBlas` hops is unchanged. + +## Stop conditions + +- Research BLOCK on this spec. +- Any attempt to transplant dirty-lab hipBLASLt / FP8 Lt / layer-split with this slot change. +- GPU smoke by lab without coord ownership. +- Landing this row on a shared branch with #838/#839. + +## Evidence + +Bus: `82b2`, `713f`, `9772`, `25c9`, `5071`, `64cb`. Donor bytes hashed in `.agents/evidence/rocm-gemma4-getblas/MANIFEST.md`. From 61e32753229768c78838edf60348c68373dee04e Mon Sep 17 00:00:00 2001 From: Don Mirror Date: Fri, 14 Aug 2026 20:41:24 -0500 Subject: [PATCH 2/7] feat(rocm): #837 GetBlas dual-slot TLS + host lifetime seam MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace single-TLS destroy-on-hop with tls_slots[2] (device 1 → slot 1). Host recorder table covers first-use / 0→1→0 / stream rebind / capture. RED mutants: swapped selector, destroy-on-hop, missing SetStream, capture SetDevice. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes] --- .../CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md | 2 +- docs/FEATURES.md | 2 +- docs/USAGE.md | 6 +- include/vt/rocm/rocm_getblas_dualslot.h | 58 ++++ src/vt/rocm/rocm_matmul_hipblaslt.hip | 55 ++-- tests/CMakeLists.txt | 3 + tests/vt/test_ops_getblas_dualslot.cpp | 301 ++++++++++++++++++ 7 files changed, 397 insertions(+), 30 deletions(-) create mode 100644 include/vt/rocm/rocm_getblas_dualslot.h create mode 100644 tests/vt/test_ops_getblas_dualslot.cpp diff --git a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md index 859950519..f806e440e 100644 --- a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md +++ b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `SPIKE` spec-only | 2026-08-14 — `64cb` split repair; donor slice hashed | +| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — dual-slot GetBlas + host lifetime seam | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index ebec16432..a6c7f2688 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -358,7 +358,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | -| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 resident experts, SharedK-WMMA prefill (RDNA4); decode-graph and forward extract deferred. Env `VT_GEMMA4_*`/`VT_ATTN_*`, seam `test_gemma4_rocm_fp8_seams`. [spec](../.agents/specs/gemma4-rocm-fp8-moe.md) | +| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 experts + SharedK-WMMA (RDNA4); #837 GetBlas `tls_slots[2]` host lifetime seam. Decode-graph deferred | ## How to read this page diff --git a/docs/USAGE.md b/docs/USAGE.md index ab4f6e450..4230d26c5 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -3402,8 +3402,10 @@ interval does. Dual-GPU resident FP8 MoE and SharedK-WMMA prefill are controlled via ENVIRONMENT.md (`VT_GEMMA4_RESIDENT_*`, `VT_ATTN_*`). Defaults stay safe off RDNA4. -This PR does **not** restructure the Gemma-4 layer loop or enable decode hipGraph -(those stay lab-only until a CUDA token-exact gate can land them). +GetBlas keeps two per-thread hipBLAS handles (`tls_slots[2]`, device 1 → slot 1) +so a 0→1 hop does not destroy GPU0's handle. No new env. This PR does **not** +restructure the Gemma-4 layer loop or enable decode hipGraph (those stay lab-only +until a CUDA token-exact gate can land them). ## LTX-2.5 text conditioning diff --git a/include/vt/rocm/rocm_getblas_dualslot.h b/include/vt/rocm/rocm_getblas_dualslot.h new file mode 100644 index 000000000..f0168ba38 --- /dev/null +++ b/include/vt/rocm/rocm_getblas_dualslot.h @@ -0,0 +1,58 @@ +// #837 GetBlas dual-slot TLS — slot index + hookable lifetime engine. +// Product GetBlas in rocm_matmul_hipblaslt.hip instantiates this with HIP hooks. +// Host tests instantiate the same engine with a recorder (no GPU). +#pragma once + +namespace vt::rocm { + +// Device 1 owns slot 1. Every other device id (0, 2, ...) shares slot 0. +// Two-GPU lab assumption; do not invent a map (spec #837). +inline int GetBlasSlotIndex(int device) { return (device == 1) ? 1 : 0; } + +// Dual-slot GetBlas lifetime. Hooks must provide: +// handle_t, stream_t +// NullHandle(), NullStream(), IsNull(handle) +// StreamIsCapturing(stream) -> bool +// GetDevice() -> int +// SetDevice(int) +// Create() -> handle_t +// Destroy(handle_t) +// SetStream(handle_t, stream_t) +template +struct GetBlasDualSlotEngine { + using handle_t = typename Hooks::handle_t; + using stream_t = typename Hooks::stream_t; + + struct Tls { + int dev = -1; + stream_t stream{}; + handle_t handle{}; + }; + + Tls tls_slots[2]{}; + + handle_t Get(int device, stream_t stream, Hooks& hooks) { + Tls& tls = tls_slots[SlotIndex(device)]; + if (!hooks.StreamIsCapturing(stream)) { + const int cur = hooks.GetDevice(); + if (cur != device) hooks.SetDevice(device); + } + if (hooks.IsNull(tls.handle) || tls.dev != device) { + if (!hooks.IsNull(tls.handle)) { + hooks.Destroy(tls.handle); + tls.handle = hooks.NullHandle(); + } + if (!hooks.StreamIsCapturing(stream)) hooks.SetDevice(device); + tls.handle = hooks.Create(); + tls.dev = device; + tls.stream = hooks.NullStream(); + } + if (tls.stream != stream) { + hooks.SetStream(tls.handle, stream); + tls.stream = stream; + } + return tls.handle; + } +}; + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index 43919747b..43dd17a8b 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -27,6 +27,7 @@ #include "vt/ops.h" #include "vt/rocm/rocm_device_bind.h" +#include "vt/rocm/rocm_getblas_dualslot.h" namespace vt::rocm { namespace { @@ -67,36 +68,38 @@ struct BlasCtx { hipblasHandle_t handle = nullptr; }; -// Per-thread handle: hipBLAS is not free-threaded; decode is single-threaded per -// engine worker. Avoid global mutex + SetStream on every GEMM (100s×/token). -BlasCtx GetBlas(int device, hipStream_t stream) { - struct Tls { - int dev = -1; - hipStream_t stream = nullptr; - hipblasHandle_t handle = nullptr; - }; - static thread_local Tls tls; - // Bind current device outside capture so create/setStream see the right GPU - // after peer-MoE left current on the expert device. - if (!StreamIsCapturing(stream)) { +// Per-thread handles: hipBLAS is not free-threaded; decode is single-threaded +// per engine worker. Avoid global mutex + SetStream on every GEMM (100s×/token). +// Dual-slot (#837): hop 0→1 must not hipblasDestroy GPU0 while GPU0 GEMMs may +// still be queued. Device 1 → slot 1; every other id shares slot 0. +struct HipBlasHooks { + using handle_t = hipblasHandle_t; + using stream_t = hipStream_t; + static handle_t NullHandle() { return nullptr; } + static stream_t NullStream() { return nullptr; } + static bool IsNull(handle_t h) { return h == nullptr; } + static bool StreamIsCapturing(stream_t s) { return vt::rocm::StreamIsCapturing(s); } + static int GetDevice() { int cur = -1; - if (hipGetDevice(&cur) != hipSuccess || cur != device) (void)hipSetDevice(device); + if (hipGetDevice(&cur) != hipSuccess) return -1; + return cur; } - if (tls.handle == nullptr || tls.dev != device) { - if (tls.handle) { - (void)hipblasDestroy(tls.handle); - tls.handle = nullptr; - } - if (!StreamIsCapturing(stream)) (void)hipSetDevice(device); - CheckBlas(hipblasCreate(&tls.handle), "hipblasCreate"); - tls.dev = device; - tls.stream = nullptr; + static void SetDevice(int d) { (void)hipSetDevice(d); } + static handle_t Create() { + handle_t h = nullptr; + CheckBlas(hipblasCreate(&h), "hipblasCreate"); + return h; } - if (tls.stream != stream) { - CheckBlas(hipblasSetStream(tls.handle, stream), "hipblasSetStream"); - tls.stream = stream; + static void Destroy(handle_t h) { (void)hipblasDestroy(h); } + static void SetStream(handle_t h, stream_t s) { + CheckBlas(hipblasSetStream(h, s), "hipblasSetStream"); } - return BlasCtx{tls.handle}; +}; + +BlasCtx GetBlas(int device, hipStream_t stream) { + static thread_local GetBlasDualSlotEngine tls_slots; + HipBlasHooks hooks; + return BlasCtx{tls_slots.Get(device, stream, hooks)}; } std::string ComboName(const Tensor& a, const Tensor& b, const Tensor& out) { diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 4a94639d9..6c7c47c9d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1775,6 +1775,9 @@ vllm_cpp_add_test(test_ops_paged_attn vt/test_ops_paged_attn.cpp) vllm_cpp_add_test(test_ops_paged_attn_dtype vt/test_ops_paged_attn_dtype.cpp) vllm_cpp_add_test(test_ops_sample vt/test_ops_sample.cpp) vllm_cpp_add_test(test_gemma4_rocm_fp8_seams vt/test_gemma4_rocm_fp8_seams.cpp) +vllm_cpp_add_test(test_ops_getblas_dualslot vt/test_ops_getblas_dualslot.cpp) +target_compile_definitions(test_ops_getblas_dualslot PRIVATE + VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") vllm_cpp_add_test(test_ops_penalties vt/test_ops_penalties.cpp) vllm_cpp_add_test(test_npy parity/test_npy.cpp) diff --git a/tests/vt/test_ops_getblas_dualslot.cpp b/tests/vt/test_ops_getblas_dualslot.cpp new file mode 100644 index 000000000..7f3c6aec9 --- /dev/null +++ b/tests/vt/test_ops_getblas_dualslot.cpp @@ -0,0 +1,301 @@ +// #837 GetBlas dual-slot host lifetime seam. No GPU required. +// Product engine (GetBlasDualSlotEngine + GetBlasSlotIndex) vs mutants that +// must violate the same table (research 64cb stop-ship 5). +#include +#include +#include +#include +#include + +#include + +#include "vt/rocm/rocm_getblas_dualslot.h" + +#ifndef VLLM_CPP_SOURCE_DIR +#define VLLM_CPP_SOURCE_DIR "." +#endif + +namespace { + +std::string ReadText(const char* rel) { + const std::string path = std::string(VLLM_CPP_SOURCE_DIR) + "/" + rel; + std::ifstream in(path); + REQUIRE(in.good()); + std::ostringstream ss; + ss << in.rdbuf(); + return ss.str(); +} + +struct Rec { + int create[2] = {0, 0}; + int destroy[2] = {0, 0}; + int set_stream[2] = {0, 0}; + int set_device_n = 0; + int last_set_device = -999; + uintptr_t handle[2] = {0, 0}; + int stream[2] = {0, 0}; + int next_id = 1; + bool capturing = false; + int cur_dev = -1; +}; + +struct FakeHooks { + using handle_t = uintptr_t; + using stream_t = int; + Rec* rec = nullptr; + int last_create_slot = -1; + + static handle_t NullHandle() { return 0; } + static stream_t NullStream() { return 0; } + static bool IsNull(handle_t h) { return h == 0; } + + bool StreamIsCapturing(stream_t) const { return rec->capturing; } + int GetDevice() const { return rec->cur_dev; } + void SetDevice(int d) { + rec->set_device_n++; + rec->last_set_device = d; + rec->cur_dev = d; + } + handle_t Create() { + const int slot = last_create_slot >= 0 ? last_create_slot : 0; + rec->create[slot]++; + const handle_t h = static_cast(rec->next_id++); + rec->handle[slot] = h; + return h; + } + void Destroy(handle_t h) { + for (int i = 0; i < 2; ++i) { + if (rec->handle[i] == h) { + rec->destroy[i]++; + rec->handle[i] = 0; + return; + } + } + rec->destroy[0]++; + } + void SetStream(handle_t h, stream_t s) { + for (int i = 0; i < 2; ++i) { + if (rec->handle[i] == h) { + rec->set_stream[i]++; + rec->stream[i] = s; + return; + } + } + } +}; + +// Tag Create with the slot the engine is about to fill. The engine calls +// Create after selecting the slot; we infer slot from which handle is still 0 +// and which device is current — set by a thin wrapper. +struct TrackingHooks : FakeHooks { + handle_t Create() { + // Slot = last SetDevice target when it is 0 or 1; else 0 (devices ≥2). + last_create_slot = (rec->last_set_device == 1) ? 1 : 0; + return FakeHooks::Create(); + } +}; + +inline int SwappedSlotIndex(int device) { return (device == 0) ? 1 : 0; } + +struct Table { + Rec rec; + uintptr_t h0 = 0; + uintptr_t h1 = 0; +}; + +template +void RunHops(Eng& eng, Hooks& hooks, Table* t, bool change_s0_stream) { + hooks.rec = &t->rec; + t->h0 = eng.Get(0, /*s0=*/1, hooks); + t->h1 = eng.Get(1, /*s1=*/2, hooks); + t->h0 = eng.Get(0, /*s0=*/1, hooks); // second visit; product keeps identity + t->h1 = eng.Get(1, /*s1=*/2, hooks); + if (change_s0_stream) (void)eng.Get(0, /*s0'=*/3, hooks); +} + +bool ProductTableHolds(const Table& t, bool after_stream_change) { + if (t.rec.create[0] != 1) return false; + if (t.rec.create[1] != 1) return false; + if (t.rec.destroy[0] != 0) return false; + if (t.rec.destroy[1] != 0) return false; + if (t.rec.handle[0] != t.h0) return false; + if (t.rec.handle[1] != t.h1) return false; + if (t.h0 == 0 || t.h1 == 0 || t.h0 == t.h1) return false; + if (after_stream_change) { + if (t.rec.stream[0] != 3) return false; + if (t.rec.stream[1] != 2) return false; + } + return true; +} + +// Old single-TLS: destroy on device hop. +struct SingleTlsEngine { + using handle_t = uintptr_t; + using stream_t = int; + struct Tls { + int dev = -1; + stream_t stream = 0; + handle_t handle = 0; + } tls; + handle_t Get(int device, stream_t stream, TrackingHooks& hooks) { + if (!hooks.StreamIsCapturing(stream)) { + if (hooks.GetDevice() != device) hooks.SetDevice(device); + } + if (hooks.IsNull(tls.handle) || tls.dev != device) { + if (!hooks.IsNull(tls.handle)) { + hooks.Destroy(tls.handle); + tls.handle = 0; + } + if (!hooks.StreamIsCapturing(stream)) hooks.SetDevice(device); + hooks.last_create_slot = (device == 1) ? 1 : 0; + tls.handle = hooks.Create(); + tls.dev = device; + tls.stream = 0; + } + if (tls.stream != stream) { + hooks.SetStream(tls.handle, stream); + tls.stream = stream; + } + return tls.handle; + } +}; + +struct NoStreamHooks : TrackingHooks { + void SetStream(handle_t, stream_t) {} +}; + +struct CaptureAlwaysSetEngine { + vt::rocm::GetBlasDualSlotEngine inner; + uintptr_t Get(int device, int stream, TrackingHooks& hooks) { + // Mutant: ignore capture, always SetDevice. + hooks.SetDevice(device); + return inner.Get(device, stream, hooks); + } +}; + +} // namespace + +TEST_CASE("getblas slot index") { + CHECK(vt::rocm::GetBlasSlotIndex(0) == 0); + CHECK(vt::rocm::GetBlasSlotIndex(1) == 1); + CHECK(vt::rocm::GetBlasSlotIndex(2) == 0); + CHECK(vt::rocm::GetBlasSlotIndex(-1) == 0); +} + +TEST_CASE("getblas first use fills slot 0 only") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Rec rec; + hooks.rec = &rec; + const auto h0 = eng.Get(0, 1, hooks); + CHECK(h0 != 0); + CHECK(eng.tls_slots[0].handle == h0); + CHECK(eng.tls_slots[1].handle == 0); + CHECK(rec.create[0] == 1); + CHECK(rec.create[1] == 0); +} + +TEST_CASE("getblas product lifetime table") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Table t; + RunHops(eng, hooks, &t, /*change_s0_stream=*/true); + CHECK(ProductTableHolds(t, /*after_stream_change=*/true)); + CHECK(eng.tls_slots[0].handle == t.h0); + CHECK(eng.tls_slots[1].handle == t.h1); + CHECK(t.rec.create[0] == 1); + CHECK(t.rec.create[1] == 1); + CHECK(t.rec.destroy[0] == 0); + CHECK(t.rec.destroy[1] == 0); + CHECK(t.rec.stream[0] == 3); + CHECK(t.rec.stream[1] == 2); +} + +TEST_CASE("getblas hop 1->0->1 keeps both handles") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Rec rec; + hooks.rec = &rec; + const auto h0 = eng.Get(0, 1, hooks); + const auto h1 = eng.Get(1, 2, hooks); + CHECK(eng.Get(0, 1, hooks) == h0); + CHECK(eng.Get(1, 2, hooks) == h1); + CHECK(rec.destroy[0] == 0); + CHECK(rec.destroy[1] == 0); + CHECK(rec.create[0] == 1); + CHECK(rec.create[1] == 1); +} + +TEST_CASE("getblas capture path does not SetDevice") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Rec rec; + rec.capturing = true; + rec.cur_dev = 0; + hooks.rec = &rec; + (void)eng.Get(1, 9, hooks); + CHECK(rec.set_device_n == 0); + CHECK(rec.last_set_device == -999); +} + +TEST_CASE("getblas RED swapped selector fills the wrong slot") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Rec rec; + hooks.rec = &rec; + const auto h = eng.Get(0, 1, hooks); + CHECK(eng.tls_slots[0].handle == 0); + CHECK(eng.tls_slots[1].handle == h); +} + +TEST_CASE("getblas RED destroy-on-hop fails table") { + SingleTlsEngine eng; + TrackingHooks hooks; + Table t; + RunHops(eng, hooks, &t, /*change_s0_stream=*/false); + CHECK(t.rec.destroy[0] >= 1); + CHECK_FALSE(ProductTableHolds(t, /*after_stream_change=*/false)); +} + +TEST_CASE("getblas RED missing stream rebind fails table") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + NoStreamHooks hooks; + Table t; + RunHops(eng, hooks, &t, /*change_s0_stream=*/true); + CHECK(t.rec.stream[0] != 3); + CHECK_FALSE(ProductTableHolds(t, /*after_stream_change=*/true)); +} + +TEST_CASE("getblas RED capture SetDevice fails capture invariant") { + CaptureAlwaysSetEngine eng; + TrackingHooks hooks; + Rec rec; + rec.capturing = true; + rec.cur_dev = 0; + hooks.rec = &rec; + (void)eng.Get(1, 9, hooks); + CHECK(rec.set_device_n >= 1); +} + +TEST_CASE("getblas product source uses dual-slot engine") { + const std::string hip = ReadText("src/vt/rocm/rocm_matmul_hipblaslt.hip"); + CHECK(hip.find("GetBlasDualSlotEngine") != std::string::npos); + CHECK(hip.find("tls_slots") != std::string::npos); + CHECK(hip.find("rocm_getblas_dualslot.h") != std::string::npos); + CHECK(hip.find("static thread_local Tls tls;") == std::string::npos); +} + +TEST_CASE("getblas GPU probe skipped without HIP_VISIBLE_DEVICES") { + const char* env = std::getenv("HIP_VISIBLE_DEVICES"); + if (env == nullptr || env[0] == '\0') return; +#if !defined(VLLM_CPP_HIP) + // Host binary cannot call product GetBlas. Coord owns the live gfx1201 probe. + return; +#endif +} From 246b2adc3711a39355bdd8aaa3e760d0ca8a479f Mon Sep 17 00:00:00 2001 From: Don Mirror Date: Fri, 14 Aug 2026 21:09:50 -0500 Subject: [PATCH 3/7] =?UTF-8?q?fix(rocm):=20#837=20product-call=20seam=20+?= =?UTF-8?q?=200=E2=86=921=E2=86=920=E2=86=921=20identity=20probe?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fold research c24b: GetBlas and host tests both execute RocmProductGetBlasOn (device+stream forwarded). RED mutants: forward device 0, forward null stream. Host 0→1→0→1 always runs; live hipblas probe compiles under VLLM_CPP_HIP. 3c08ce35 is not a review target. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes] --- .../CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md | 2 +- docs/FEATURES.md | 2 +- docs/USAGE.md | 3 +- include/vt/rocm/rocm_getblas_dualslot.h | 21 +- src/vt/rocm/rocm_matmul_hipblaslt.hip | 3 +- tests/vt/test_ops_getblas_dualslot.cpp | 179 ++++++++++++++---- 6 files changed, 169 insertions(+), 41 deletions(-) diff --git a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md index f806e440e..7d2fe8a3e 100644 --- a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md +++ b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — dual-slot GetBlas + host lifetime seam | +| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — c24b product-call seam + 0→1→0→1 probe | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index a6c7f2688..84b18eebc 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -358,7 +358,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | -| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 experts + SharedK-WMMA (RDNA4); #837 GetBlas `tls_slots[2]` host lifetime seam. Decode-graph deferred | +| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 experts + SharedK-WMMA (RDNA4); #837 GetBlas `tls_slots[2]` + product-call seam. Decode-graph deferred | ## How to read this page diff --git a/docs/USAGE.md b/docs/USAGE.md index 4230d26c5..614f44bff 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -3403,7 +3403,8 @@ interval does. Dual-GPU resident FP8 MoE and SharedK-WMMA prefill are controlled via ENVIRONMENT.md (`VT_GEMMA4_RESIDENT_*`, `VT_ATTN_*`). Defaults stay safe off RDNA4. GetBlas keeps two per-thread hipBLAS handles (`tls_slots[2]`, device 1 → slot 1) -so a 0→1 hop does not destroy GPU0's handle. No new env. This PR does **not** +so a 0→1 hop does not destroy GPU0's handle. Production `GetBlas` and host tests +share `RocmProductGetBlasOn` (device+stream forwarded). No new env. This PR does **not** restructure the Gemma-4 layer loop or enable decode hipGraph (those stay lab-only until a CUDA token-exact gate can land them). diff --git a/include/vt/rocm/rocm_getblas_dualslot.h b/include/vt/rocm/rocm_getblas_dualslot.h index f0168ba38..8e1b8e70c 100644 --- a/include/vt/rocm/rocm_getblas_dualslot.h +++ b/include/vt/rocm/rocm_getblas_dualslot.h @@ -1,6 +1,6 @@ // #837 GetBlas dual-slot TLS — slot index + hookable lifetime engine. -// Product GetBlas in rocm_matmul_hipblaslt.hip instantiates this with HIP hooks. -// Host tests instantiate the same engine with a recorder (no GPU). +// Product GetBlas and host tests both execute RocmProductGetBlasOn +// (research c24b: substring wiring is not a load-bearing seam). #pragma once namespace vt::rocm { @@ -55,4 +55,21 @@ struct GetBlasDualSlotEngine { } }; +// Product-call seam: production GetBlas and host tests both execute this. +// Forwards device + stream unchanged. Mutating either argument is RED. +struct RocmGetBlasForward { + template + static typename Engine::handle_t apply(Engine& engine, int device, + typename Engine::stream_t stream, + Hooks& hooks) { + return engine.Get(device, stream, hooks); + } +}; + +template +inline typename Engine::handle_t RocmProductGetBlasOn( + Engine& engine, int device, typename Engine::stream_t stream, Hooks& hooks) { + return Forward::apply(engine, device, stream, hooks); +} + } // namespace vt::rocm diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index 43dd17a8b..9e6af1a3c 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -99,7 +99,8 @@ struct HipBlasHooks { BlasCtx GetBlas(int device, hipStream_t stream) { static thread_local GetBlasDualSlotEngine tls_slots; HipBlasHooks hooks; - return BlasCtx{tls_slots.Get(device, stream, hooks)}; + // Shared seam: tests execute RocmProductGetBlasOn with the same forwarding. + return BlasCtx{vt::rocm::RocmProductGetBlasOn(tls_slots, device, stream, hooks)}; } std::string ComboName(const Tensor& a, const Tensor& b, const Tensor& out) { diff --git a/tests/vt/test_ops_getblas_dualslot.cpp b/tests/vt/test_ops_getblas_dualslot.cpp index 7f3c6aec9..41e428538 100644 --- a/tests/vt/test_ops_getblas_dualslot.cpp +++ b/tests/vt/test_ops_getblas_dualslot.cpp @@ -1,6 +1,5 @@ -// #837 GetBlas dual-slot host lifetime seam. No GPU required. -// Product engine (GetBlasDualSlotEngine + GetBlasSlotIndex) vs mutants that -// must violate the same table (research 64cb stop-ship 5). +// #837 GetBlas dual-slot host lifetime + product-call seam (research c24b). +// Production GetBlas and these tests both execute RocmProductGetBlasOn. #include #include #include @@ -11,6 +10,11 @@ #include "vt/rocm/rocm_getblas_dualslot.h" +#if defined(VLLM_CPP_HIP) +#include +#include +#endif + #ifndef VLLM_CPP_SOURCE_DIR #define VLLM_CPP_SOURCE_DIR "." #endif @@ -84,12 +88,8 @@ struct FakeHooks { } }; -// Tag Create with the slot the engine is about to fill. The engine calls -// Create after selecting the slot; we infer slot from which handle is still 0 -// and which device is current — set by a thin wrapper. struct TrackingHooks : FakeHooks { handle_t Create() { - // Slot = last SetDevice target when it is 0 or 1; else 0 (devices ≥2). last_create_slot = (rec->last_set_device == 1) ? 1 : 0; return FakeHooks::Create(); } @@ -103,14 +103,17 @@ struct Table { uintptr_t h1 = 0; }; -template -void RunHops(Eng& eng, Hooks& hooks, Table* t, bool change_s0_stream) { +// Product hops go through RocmProductGetBlasOn — same function as GetBlas. +template +void RunProductHops(Eng& eng, Hooks& hooks, Table* t, bool change_s0_stream) { hooks.rec = &t->rec; - t->h0 = eng.Get(0, /*s0=*/1, hooks); - t->h1 = eng.Get(1, /*s1=*/2, hooks); - t->h0 = eng.Get(0, /*s0=*/1, hooks); // second visit; product keeps identity - t->h1 = eng.Get(1, /*s1=*/2, hooks); - if (change_s0_stream) (void)eng.Get(0, /*s0'=*/3, hooks); + t->h0 = vt::rocm::RocmProductGetBlasOn(eng, 0, /*s0=*/1, hooks); + t->h1 = vt::rocm::RocmProductGetBlasOn(eng, 1, /*s1=*/2, hooks); + t->h0 = vt::rocm::RocmProductGetBlasOn(eng, 0, /*s0=*/1, hooks); + t->h1 = vt::rocm::RocmProductGetBlasOn(eng, 1, /*s1=*/2, hooks); + if (change_s0_stream) { + (void)vt::rocm::RocmProductGetBlasOn(eng, 0, /*s0'=*/3, hooks); + } } bool ProductTableHolds(const Table& t, bool after_stream_change) { @@ -128,7 +131,6 @@ bool ProductTableHolds(const Table& t, bool after_stream_change) { return true; } -// Old single-TLS: destroy on device hop. struct SingleTlsEngine { using handle_t = uintptr_t; using stream_t = int; @@ -167,12 +169,30 @@ struct NoStreamHooks : TrackingHooks { struct CaptureAlwaysSetEngine { vt::rocm::GetBlasDualSlotEngine inner; uintptr_t Get(int device, int stream, TrackingHooks& hooks) { - // Mutant: ignore capture, always SetDevice. hooks.SetDevice(device); return inner.Get(device, stream, hooks); } }; +// c24b mutants: product call forwards the wrong device or a null stream. +struct ForwardDev0 { + template + static typename Engine::handle_t apply(Engine& engine, int /*device*/, + typename Engine::stream_t stream, + Hooks& hooks) { + return engine.Get(0, stream, hooks); + } +}; + +struct ForwardNullStream { + template + static typename Engine::handle_t apply(Engine& engine, int device, + typename Engine::stream_t /*stream*/, + Hooks& hooks) { + return engine.Get(device, Hooks::NullStream(), hooks); + } +}; + } // namespace TEST_CASE("getblas slot index") { @@ -188,7 +208,7 @@ TEST_CASE("getblas first use fills slot 0 only") { TrackingHooks hooks; Rec rec; hooks.rec = &rec; - const auto h0 = eng.Get(0, 1, hooks); + const auto h0 = vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); CHECK(h0 != 0); CHECK(eng.tls_slots[0].handle == h0); CHECK(eng.tls_slots[1].handle == 0); @@ -196,12 +216,12 @@ TEST_CASE("getblas first use fills slot 0 only") { CHECK(rec.create[1] == 0); } -TEST_CASE("getblas product lifetime table") { +TEST_CASE("getblas product lifetime table via product-call seam") { using Eng = vt::rocm::GetBlasDualSlotEngine; Eng eng; TrackingHooks hooks; Table t; - RunHops(eng, hooks, &t, /*change_s0_stream=*/true); + RunProductHops(eng, hooks, &t, /*change_s0_stream=*/true); CHECK(ProductTableHolds(t, /*after_stream_change=*/true)); CHECK(eng.tls_slots[0].handle == t.h0); CHECK(eng.tls_slots[1].handle == t.h1); @@ -219,10 +239,10 @@ TEST_CASE("getblas hop 1->0->1 keeps both handles") { TrackingHooks hooks; Rec rec; hooks.rec = &rec; - const auto h0 = eng.Get(0, 1, hooks); - const auto h1 = eng.Get(1, 2, hooks); - CHECK(eng.Get(0, 1, hooks) == h0); - CHECK(eng.Get(1, 2, hooks) == h1); + const auto h0 = vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); + const auto h1 = vt::rocm::RocmProductGetBlasOn(eng, 1, 2, hooks); + CHECK(vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks) == h0); + CHECK(vt::rocm::RocmProductGetBlasOn(eng, 1, 2, hooks) == h1); CHECK(rec.destroy[0] == 0); CHECK(rec.destroy[1] == 0); CHECK(rec.create[0] == 1); @@ -237,7 +257,7 @@ TEST_CASE("getblas capture path does not SetDevice") { rec.capturing = true; rec.cur_dev = 0; hooks.rec = &rec; - (void)eng.Get(1, 9, hooks); + (void)vt::rocm::RocmProductGetBlasOn(eng, 1, 9, hooks); CHECK(rec.set_device_n == 0); CHECK(rec.last_set_device == -999); } @@ -248,7 +268,7 @@ TEST_CASE("getblas RED swapped selector fills the wrong slot") { TrackingHooks hooks; Rec rec; hooks.rec = &rec; - const auto h = eng.Get(0, 1, hooks); + const auto h = vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); CHECK(eng.tls_slots[0].handle == 0); CHECK(eng.tls_slots[1].handle == h); } @@ -257,7 +277,7 @@ TEST_CASE("getblas RED destroy-on-hop fails table") { SingleTlsEngine eng; TrackingHooks hooks; Table t; - RunHops(eng, hooks, &t, /*change_s0_stream=*/false); + RunProductHops(eng, hooks, &t, /*change_s0_stream=*/false); CHECK(t.rec.destroy[0] >= 1); CHECK_FALSE(ProductTableHolds(t, /*after_stream_change=*/false)); } @@ -267,7 +287,7 @@ TEST_CASE("getblas RED missing stream rebind fails table") { Eng eng; NoStreamHooks hooks; Table t; - RunHops(eng, hooks, &t, /*change_s0_stream=*/true); + RunProductHops(eng, hooks, &t, /*change_s0_stream=*/true); CHECK(t.rec.stream[0] != 3); CHECK_FALSE(ProductTableHolds(t, /*after_stream_change=*/true)); } @@ -283,19 +303,108 @@ TEST_CASE("getblas RED capture SetDevice fails capture invariant") { CHECK(rec.set_device_n >= 1); } -TEST_CASE("getblas product source uses dual-slot engine") { +TEST_CASE("getblas RED product call forwards device 0") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Table t; + RunProductHops(eng, hooks, &t, false); + CHECK(t.h0 == t.h1); + CHECK_FALSE(ProductTableHolds(t, false)); +} + +TEST_CASE("getblas RED product call forwards null stream") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Table t; + RunProductHops(eng, hooks, &t, true); + CHECK(t.rec.stream[0] != 3); + CHECK_FALSE(ProductTableHolds(t, true)); +} + +TEST_CASE("getblas product source uses RocmProductGetBlasOn with device+stream") { const std::string hip = ReadText("src/vt/rocm/rocm_matmul_hipblaslt.hip"); - CHECK(hip.find("GetBlasDualSlotEngine") != std::string::npos); - CHECK(hip.find("tls_slots") != std::string::npos); - CHECK(hip.find("rocm_getblas_dualslot.h") != std::string::npos); + CHECK(hip.find("RocmProductGetBlasOn") != std::string::npos); + CHECK(hip.find("RocmProductGetBlasOn(tls_slots, device, stream, hooks)") != + std::string::npos); + CHECK(hip.find("tls_slots.Get(0,") == std::string::npos); + CHECK(hip.find("Get(device, nullptr") == std::string::npos); CHECK(hip.find("static thread_local Tls tls;") == std::string::npos); } -TEST_CASE("getblas GPU probe skipped without HIP_VISIBLE_DEVICES") { +TEST_CASE("getblas product 0-1-0-1 handle identity") { const char* env = std::getenv("HIP_VISIBLE_DEVICES"); +#if defined(VLLM_CPP_HIP) if (env == nullptr || env[0] == '\0') return; -#if !defined(VLLM_CPP_HIP) - // Host binary cannot call product GetBlas. Coord owns the live gfx1201 probe. - return; + // Live hipblas probe: same RocmProductGetBlasOn + real hooks. + // Compiled only into the HIP test binary; coord runs it on dual devices. + struct HipBlasHooks { + using handle_t = void*; + using stream_t = void*; + static handle_t NullHandle() { return nullptr; } + static stream_t NullStream() { return nullptr; } + static bool IsNull(handle_t h) { return h == nullptr; } + bool StreamIsCapturing(stream_t) const { return false; } + int GetDevice() const { + int cur = -1; + (void)hipGetDevice(&cur); + return cur; + } + void SetDevice(int d) { (void)hipSetDevice(d); } + handle_t Create() { + hipblasHandle_t h = nullptr; + if (hipblasCreate(&h) != HIPBLAS_STATUS_SUCCESS) return nullptr; + return static_cast(h); + } + void Destroy(handle_t h) { + if (h) (void)hipblasDestroy(static_cast(h)); + } + void SetStream(handle_t h, stream_t s) { + (void)hipblasSetStream(static_cast(h), + static_cast(s)); + } + }; + int ndev = 0; + REQUIRE(hipGetDeviceCount(&ndev) == hipSuccess); + REQUIRE(ndev >= 2); + hipStream_t s0 = nullptr; + hipStream_t s1 = nullptr; + REQUIRE(hipSetDevice(0) == hipSuccess); + REQUIRE(hipStreamCreate(&s0) == hipSuccess); + REQUIRE(hipSetDevice(1) == hipSuccess); + REQUIRE(hipStreamCreate(&s1) == hipSuccess); + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + HipBlasHooks hooks; + const auto h0a = vt::rocm::RocmProductGetBlasOn(eng, 0, static_cast(s0), hooks); + const auto h1a = vt::rocm::RocmProductGetBlasOn(eng, 1, static_cast(s1), hooks); + const auto h0b = vt::rocm::RocmProductGetBlasOn(eng, 0, static_cast(s0), hooks); + const auto h1b = vt::rocm::RocmProductGetBlasOn(eng, 1, static_cast(s1), hooks); + CHECK(h0a != nullptr); + CHECK(h1a != nullptr); + CHECK(h0a == h0b); + CHECK(h1a == h1b); + CHECK(h0a != h1a); + (void)hipSetDevice(0); + (void)hipStreamDestroy(s0); + (void)hipSetDevice(1); + (void)hipStreamDestroy(s1); +#else + (void)env; + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Rec rec; + hooks.rec = &rec; + const auto h0a = vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); + const auto h1a = vt::rocm::RocmProductGetBlasOn(eng, 1, 2, hooks); + const auto h0b = vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); + const auto h1b = vt::rocm::RocmProductGetBlasOn(eng, 1, 2, hooks); + CHECK(h0a == h0b); + CHECK(h1a == h1b); + CHECK(h0a != h1a); + CHECK(rec.destroy[0] == 0); + CHECK(rec.destroy[1] == 0); #endif } From 2eae0818a5d6d911651974ea9a297144c1e48ba0 Mon Sep 17 00:00:00 2001 From: Don Mirror Date: Fri, 14 Aug 2026 22:31:01 -0500 Subject: [PATCH 4/7] fix(rocm): #837 product GetBlas accessor + exit-77 HIP probe MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fold research 4d82/e819: ProductGetBlasHandle calls file-local GetBlas (real HipBlasHooks + static thread_local tls_slots). Host table asserts SetDevice revisit. HIP 0→1→0→1 is a separate CTest; missing devices exit 77. c956ba6d is not a review target. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes] --- .../CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md | 2 +- docs/FEATURES.md | 2 +- docs/USAGE.md | 5 +- include/vt/rocm/rocm_getblas_product.h | 15 ++++ src/vt/rocm/rocm_matmul_hipblaslt.hip | 14 +++- tests/CMakeLists.txt | 4 + tests/vt/test_ops_getblas_dualslot.cpp | 82 +++++-------------- tests/vt/test_ops_getblas_product.cpp | 82 +++++++++++++++++++ 8 files changed, 136 insertions(+), 70 deletions(-) create mode 100644 include/vt/rocm/rocm_getblas_product.h create mode 100644 tests/vt/test_ops_getblas_product.cpp diff --git a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md index 7d2fe8a3e..d1e5dff6e 100644 --- a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md +++ b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — c24b product-call seam + 0→1→0→1 probe | +| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — 4d82/e819 product GetBlas accessor + exit-77 HIP probe | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 84b18eebc..6454c8b47 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -358,7 +358,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | -| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 experts + SharedK-WMMA (RDNA4); #837 GetBlas `tls_slots[2]` + product-call seam. Decode-graph deferred | +| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 experts + SharedK-WMMA (RDNA4); #837 GetBlas `tls_slots[2]` + product `GetBlas` probe. Decode-graph deferred | ## How to read this page diff --git a/docs/USAGE.md b/docs/USAGE.md index 614f44bff..b825f66b6 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -3403,8 +3403,9 @@ interval does. Dual-GPU resident FP8 MoE and SharedK-WMMA prefill are controlled via ENVIRONMENT.md (`VT_GEMMA4_RESIDENT_*`, `VT_ATTN_*`). Defaults stay safe off RDNA4. GetBlas keeps two per-thread hipBLAS handles (`tls_slots[2]`, device 1 → slot 1) -so a 0→1 hop does not destroy GPU0's handle. Production `GetBlas` and host tests -share `RocmProductGetBlasOn` (device+stream forwarded). No new env. This PR does **not** +so a 0→1 hop does not destroy GPU0's handle. `ProductGetBlasHandle` is the +test accessor for that file-local `GetBlas`. HIP live probe is a separate CTest +target (exit 77 if `HIP_VISIBLE_DEVICES` empty). No new env. This PR does **not** restructure the Gemma-4 layer loop or enable decode hipGraph (those stay lab-only until a CUDA token-exact gate can land them). diff --git a/include/vt/rocm/rocm_getblas_product.h b/include/vt/rocm/rocm_getblas_product.h new file mode 100644 index 000000000..f0cf96269 --- /dev/null +++ b/include/vt/rocm/rocm_getblas_product.h @@ -0,0 +1,15 @@ +// #837 test/product accessor for file-local GetBlas. +// Defined in src/vt/rocm/rocm_matmul_hipblaslt.hip (HIP builds only). +#pragma once + +#if defined(VLLM_CPP_HIP) +#include +#include + +namespace vt::rocm { + +// Executes production GetBlas: real HipBlasHooks + static thread_local tls_slots. +hipblasHandle_t ProductGetBlasHandle(int device, hipStream_t stream); + +} // namespace vt::rocm +#endif diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index 9e6af1a3c..2873bfa7c 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -84,7 +84,7 @@ struct HipBlasHooks { if (hipGetDevice(&cur) != hipSuccess) return -1; return cur; } - static void SetDevice(int d) { (void)hipSetDevice(d); } + static void SetDevice(int d) { CheckHip(hipSetDevice(d), "hipSetDevice"); } static handle_t Create() { handle_t h = nullptr; CheckBlas(hipblasCreate(&h), "hipblasCreate"); @@ -99,10 +99,18 @@ struct HipBlasHooks { BlasCtx GetBlas(int device, hipStream_t stream) { static thread_local GetBlasDualSlotEngine tls_slots; HipBlasHooks hooks; - // Shared seam: tests execute RocmProductGetBlasOn with the same forwarding. - return BlasCtx{vt::rocm::RocmProductGetBlasOn(tls_slots, device, stream, hooks)}; + return BlasCtx{RocmProductGetBlasOn(tls_slots, device, stream, hooks)}; } +} // namespace + +hipblasHandle_t ProductGetBlasHandle(int device, hipStream_t stream) { + // Same-TU call into file-local GetBlas (HipBlasHooks + tls_slots). + return GetBlas(device, stream).handle; +} + +namespace { + std::string ComboName(const Tensor& a, const Tensor& b, const Tensor& out) { return std::string("(") + Name(a.dtype) + "," + Name(b.dtype) + ")->" + Name(out.dtype); } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 6c7c47c9d..f880473b2 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1778,6 +1778,10 @@ vllm_cpp_add_test(test_gemma4_rocm_fp8_seams vt/test_gemma4_rocm_fp8_seams.cpp) vllm_cpp_add_test(test_ops_getblas_dualslot vt/test_ops_getblas_dualslot.cpp) target_compile_definitions(test_ops_getblas_dualslot PRIVATE VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") +if(VLLM_CPP_HIP) + # Separate HIP target: calls production GetBlas. Missing devices → exit 77. + vllm_cpp_add_test(test_ops_getblas_product vt/test_ops_getblas_product.cpp) +endif() vllm_cpp_add_test(test_ops_penalties vt/test_ops_penalties.cpp) vllm_cpp_add_test(test_npy parity/test_npy.cpp) diff --git a/tests/vt/test_ops_getblas_dualslot.cpp b/tests/vt/test_ops_getblas_dualslot.cpp index 41e428538..8c7cf55dd 100644 --- a/tests/vt/test_ops_getblas_dualslot.cpp +++ b/tests/vt/test_ops_getblas_dualslot.cpp @@ -10,10 +10,6 @@ #include "vt/rocm/rocm_getblas_dualslot.h" -#if defined(VLLM_CPP_HIP) -#include -#include -#endif #ifndef VLLM_CPP_SOURCE_DIR #define VLLM_CPP_SOURCE_DIR "." @@ -333,65 +329,25 @@ TEST_CASE("getblas product source uses RocmProductGetBlasOn with device+stream") CHECK(hip.find("static thread_local Tls tls;") == std::string::npos); } -TEST_CASE("getblas product 0-1-0-1 handle identity") { - const char* env = std::getenv("HIP_VISIBLE_DEVICES"); -#if defined(VLLM_CPP_HIP) - if (env == nullptr || env[0] == '\0') return; - // Live hipblas probe: same RocmProductGetBlasOn + real hooks. - // Compiled only into the HIP test binary; coord runs it on dual devices. - struct HipBlasHooks { - using handle_t = void*; - using stream_t = void*; - static handle_t NullHandle() { return nullptr; } - static stream_t NullStream() { return nullptr; } - static bool IsNull(handle_t h) { return h == nullptr; } - bool StreamIsCapturing(stream_t) const { return false; } - int GetDevice() const { - int cur = -1; - (void)hipGetDevice(&cur); - return cur; - } - void SetDevice(int d) { (void)hipSetDevice(d); } - handle_t Create() { - hipblasHandle_t h = nullptr; - if (hipblasCreate(&h) != HIPBLAS_STATUS_SUCCESS) return nullptr; - return static_cast(h); - } - void Destroy(handle_t h) { - if (h) (void)hipblasDestroy(static_cast(h)); - } - void SetStream(handle_t h, stream_t s) { - (void)hipblasSetStream(static_cast(h), - static_cast(s)); - } - }; - int ndev = 0; - REQUIRE(hipGetDeviceCount(&ndev) == hipSuccess); - REQUIRE(ndev >= 2); - hipStream_t s0 = nullptr; - hipStream_t s1 = nullptr; - REQUIRE(hipSetDevice(0) == hipSuccess); - REQUIRE(hipStreamCreate(&s0) == hipSuccess); - REQUIRE(hipSetDevice(1) == hipSuccess); - REQUIRE(hipStreamCreate(&s1) == hipSuccess); - using Eng = vt::rocm::GetBlasDualSlotEngine; +TEST_CASE("getblas hop 0-1-0-1 restores device on revisit") { + using Eng = vt::rocm::GetBlasDualSlotEngine; Eng eng; - HipBlasHooks hooks; - const auto h0a = vt::rocm::RocmProductGetBlasOn(eng, 0, static_cast(s0), hooks); - const auto h1a = vt::rocm::RocmProductGetBlasOn(eng, 1, static_cast(s1), hooks); - const auto h0b = vt::rocm::RocmProductGetBlasOn(eng, 0, static_cast(s0), hooks); - const auto h1b = vt::rocm::RocmProductGetBlasOn(eng, 1, static_cast(s1), hooks); - CHECK(h0a != nullptr); - CHECK(h1a != nullptr); - CHECK(h0a == h0b); - CHECK(h1a == h1b); - CHECK(h0a != h1a); - (void)hipSetDevice(0); - (void)hipStreamDestroy(s0); - (void)hipSetDevice(1); - (void)hipStreamDestroy(s1); -#else - (void)env; + TrackingHooks hooks; + Rec rec; + hooks.rec = &rec; + (void)vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); + CHECK(rec.cur_dev == 0); + (void)vt::rocm::RocmProductGetBlasOn(eng, 1, 2, hooks); + CHECK(rec.cur_dev == 1); + (void)vt::rocm::RocmProductGetBlasOn(eng, 0, 1, hooks); + CHECK(rec.cur_dev == 0); + CHECK(rec.last_set_device == 0); + (void)vt::rocm::RocmProductGetBlasOn(eng, 1, 2, hooks); + CHECK(rec.cur_dev == 1); + CHECK(rec.last_set_device == 1); +} + +TEST_CASE("getblas product 0-1-0-1 handle identity") { using Eng = vt::rocm::GetBlasDualSlotEngine; Eng eng; TrackingHooks hooks; @@ -406,5 +362,5 @@ TEST_CASE("getblas product 0-1-0-1 handle identity") { CHECK(h0a != h1a); CHECK(rec.destroy[0] == 0); CHECK(rec.destroy[1] == 0); -#endif + CHECK(rec.cur_dev == 1); } diff --git a/tests/vt/test_ops_getblas_product.cpp b/tests/vt/test_ops_getblas_product.cpp new file mode 100644 index 000000000..6938c0535 --- /dev/null +++ b/tests/vt/test_ops_getblas_product.cpp @@ -0,0 +1,82 @@ +// #837 live product GetBlas probe (research 4d82/e819). +// Calls vt::rocm::ProductGetBlasHandle → file-local GetBlas. +// Missing HIP_VISIBLE_DEVICES or <2 devices: exit 77 (CTest SKIP), never SUCCESS. +#include +#include + +#include +#include +#include + +#include "vt/rocm/rocm_getblas_product.h" + +namespace { + +[[noreturn]] void SkipNotRun(const char* why) { + std::fprintf(stderr, + "\n*** GATE NOT RUN — SKIPPED (exit 77), this is NOT a pass ***\n" + "%s\n", + why); + std::exit(77); +} + +void RequireHip(hipError_t err, const char* what) { + REQUIRE(err == hipSuccess); + (void)what; +} + +} // namespace + +TEST_CASE("getblas product GetBlas 0-1-0-1 handle identity") { + const char* env = std::getenv("HIP_VISIBLE_DEVICES"); + if (env == nullptr || env[0] == '\0') { + SkipNotRun("HIP_VISIBLE_DEVICES empty"); + } + int ndev = 0; + if (hipGetDeviceCount(&ndev) != hipSuccess || ndev < 2) { + SkipNotRun("need >= 2 HIP devices"); + } + + hipStream_t s0 = nullptr; + hipStream_t s1 = nullptr; + RequireHip(hipSetDevice(0), "set0"); + RequireHip(hipStreamCreate(&s0), "s0"); + RequireHip(hipSetDevice(1), "set1"); + RequireHip(hipStreamCreate(&s1), "s1"); + + const hipblasHandle_t h0a = vt::rocm::ProductGetBlasHandle(0, s0); + int cur = -1; + RequireHip(hipGetDevice(&cur), "get after 0"); + REQUIRE(h0a != nullptr); + CHECK(cur == 0); + hipStream_t bound = nullptr; + REQUIRE(hipblasGetStream(h0a, &bound) == HIPBLAS_STATUS_SUCCESS); + CHECK(bound == s0); + + const hipblasHandle_t h1a = vt::rocm::ProductGetBlasHandle(1, s1); + RequireHip(hipGetDevice(&cur), "get after 1"); + REQUIRE(h1a != nullptr); + CHECK(h1a != h0a); + CHECK(cur == 1); + REQUIRE(hipblasGetStream(h1a, &bound) == HIPBLAS_STATUS_SUCCESS); + CHECK(bound == s1); + + const hipblasHandle_t h0b = vt::rocm::ProductGetBlasHandle(0, s0); + RequireHip(hipGetDevice(&cur), "get after 0 revisit"); + CHECK(h0b == h0a); + CHECK(cur == 0); + REQUIRE(hipblasGetStream(h0b, &bound) == HIPBLAS_STATUS_SUCCESS); + CHECK(bound == s0); + + const hipblasHandle_t h1b = vt::rocm::ProductGetBlasHandle(1, s1); + RequireHip(hipGetDevice(&cur), "get after 1 revisit"); + CHECK(h1b == h1a); + CHECK(cur == 1); + REQUIRE(hipblasGetStream(h1b, &bound) == HIPBLAS_STATUS_SUCCESS); + CHECK(bound == s1); + + RequireHip(hipSetDevice(0), "cleanup0"); + RequireHip(hipStreamDestroy(s0), "ds0"); + RequireHip(hipSetDevice(1), "cleanup1"); + RequireHip(hipStreamDestroy(s1), "ds1"); +} From 522c2ce42108a582e8b10ce616ae7f5c6dd6ea16 Mon Sep 17 00:00:00 2001 From: Don Mirror Date: Fri, 14 Aug 2026 22:57:34 -0500 Subject: [PATCH 5/7] fix(rocm): #837 production StreamIsCapturing is load-bearing Fold research 6195: ProductGetBlasStreamIsCapturing calls the exact HipBlasHooks hook. HIP product probe begins capture and asserts true; always-false hook mutation is RED. Host fake-capture case unchanged. 1a1153d6 is not a review target. Adjacent #785/#523/#509/#834 noted in spec; no pickup. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes] --- .../CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md | 2 +- .agents/specs/rocm-gemma4-getblas-dualslot.md | 8 +++++ docs/USAGE.md | 2 +- include/vt/rocm/rocm_getblas_product.h | 3 ++ src/vt/rocm/rocm_matmul_hipblaslt.hip | 5 +++ tests/vt/test_ops_getblas_product.cpp | 32 +++++++++++++++++++ 6 files changed, 50 insertions(+), 2 deletions(-) diff --git a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md index d1e5dff6e..7fcbb8a4c 100644 --- a/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md +++ b/.agents/claims/CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT.md @@ -2,4 +2,4 @@ | Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update | |---|---|---|---|---|---|---|---| -| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — 4d82/e819 product GetBlas accessor + exit-77 HIP probe | +| `CLAIM-ROCM-GEMMA4-GETBLAS-DUALSLOT` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-GETBLAS-DUALSLOT`, issue #837) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-getblas` | `row/ROCM-GEMMA4-GETBLAS-DUALSLOT` | Owns ONLY: `GetBlas` `tls_slots[2]` in `src/vt/rocm/rocm_matmul_hipblaslt.hip` plus host lifetime seam tests. **EXCLUDED:** Launch/Finish (#839), indexed T (#838), #697 / `rocm_paged_attn.hip`. Independent history from the abandoned combined branch `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — 6195 production capture hook load-bearing | diff --git a/.agents/specs/rocm-gemma4-getblas-dualslot.md b/.agents/specs/rocm-gemma4-getblas-dualslot.md index c676493af..2526b020f 100644 --- a/.agents/specs/rocm-gemma4-getblas-dualslot.md +++ b/.agents/specs/rocm-gemma4-getblas-dualslot.md @@ -51,6 +51,14 @@ Keep the existing capture/`hipGetDevice`/`hipSetDevice`/`hipblasSetStream` body - Diagnostic `STAGE_SYNC` / `PREFILL_TRACE`. - More than two slots; devices ≥2. + +## Adjacent upstream (not this row) + +- **#785** (joral, OPEN issue): host-dead `VT_ROCWMMA_OK` around SharedK launch. Hard same-hunk landing-order overlap with **#697**, not with GetBlas. Do not patch that guard here. +- **#523 / #509** (VikashLoomba, OPEN drafts): custom keep-quant grouped GEMM + `rocm_moe_chain.hip`. Path intersection with this row is `docs/FEATURES.md` + `docs/USAGE.md` only. Their expert GEMM does not call `GetBlas` / `ProductGetBlasHandle`. +- **#834** (unowned): router-lookahead prefetch. Adjacent cache policy only; not TLS lifetime. + + ## Design 1. Two process-lifetime per-thread handles. Hop 0→1 must not destroy GPU0's handle while GPU0 GEMMs may still be queued (donor comment at `:74-75`). diff --git a/docs/USAGE.md b/docs/USAGE.md index b825f66b6..81b6c4385 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -3405,7 +3405,7 @@ ENVIRONMENT.md (`VT_GEMMA4_RESIDENT_*`, `VT_ATTN_*`). Defaults stay safe off RDN GetBlas keeps two per-thread hipBLAS handles (`tls_slots[2]`, device 1 → slot 1) so a 0→1 hop does not destroy GPU0's handle. `ProductGetBlasHandle` is the test accessor for that file-local `GetBlas`. HIP live probe is a separate CTest -target (exit 77 if `HIP_VISIBLE_DEVICES` empty). No new env. This PR does **not** +target (exit 77 if `HIP_VISIBLE_DEVICES` empty); it enters capture so production `StreamIsCapturing` is load-bearing. No new env. This PR does **not** restructure the Gemma-4 layer loop or enable decode hipGraph (those stay lab-only until a CUDA token-exact gate can land them). diff --git a/include/vt/rocm/rocm_getblas_product.h b/include/vt/rocm/rocm_getblas_product.h index f0cf96269..f36ce8aa8 100644 --- a/include/vt/rocm/rocm_getblas_product.h +++ b/include/vt/rocm/rocm_getblas_product.h @@ -11,5 +11,8 @@ namespace vt::rocm { // Executes production GetBlas: real HipBlasHooks + static thread_local tls_slots. hipblasHandle_t ProductGetBlasHandle(int device, hipStream_t stream); +// Exact HipBlasHooks::StreamIsCapturing (not a parallel hipStreamIsCapturing). +bool ProductGetBlasStreamIsCapturing(hipStream_t stream); + } // namespace vt::rocm #endif diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index 2873bfa7c..a1d043702 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -109,6 +109,11 @@ hipblasHandle_t ProductGetBlasHandle(int device, hipStream_t stream) { return GetBlas(device, stream).handle; } +bool ProductGetBlasStreamIsCapturing(hipStream_t stream) { + // Exact production hook. Always-false mutation must turn the HIP probe RED. + return HipBlasHooks::StreamIsCapturing(stream); +} + namespace { std::string ComboName(const Tensor& a, const Tensor& b, const Tensor& out) { diff --git a/tests/vt/test_ops_getblas_product.cpp b/tests/vt/test_ops_getblas_product.cpp index 6938c0535..5cdec3386 100644 --- a/tests/vt/test_ops_getblas_product.cpp +++ b/tests/vt/test_ops_getblas_product.cpp @@ -80,3 +80,35 @@ TEST_CASE("getblas product GetBlas 0-1-0-1 handle identity") { RequireHip(hipSetDevice(1), "cleanup1"); RequireHip(hipStreamDestroy(s1), "ds1"); } + +TEST_CASE("getblas product capture hook is load-bearing") { + const char* env = std::getenv("HIP_VISIBLE_DEVICES"); + if (env == nullptr || env[0] == '\0') { + SkipNotRun("HIP_VISIBLE_DEVICES empty"); + } + int ndev = 0; + if (hipGetDeviceCount(&ndev) != hipSuccess || ndev < 1) { + SkipNotRun("need >= 1 HIP device"); + } + + hipStream_t s0 = nullptr; + RequireHip(hipSetDevice(0), "set0"); + RequireHip(hipStreamCreate(&s0), "s0"); + const hipblasHandle_t h0 = vt::rocm::ProductGetBlasHandle(0, s0); + REQUIRE(h0 != nullptr); + CHECK(vt::rocm::ProductGetBlasStreamIsCapturing(s0) == false); + + RequireHip(hipStreamBeginCapture(s0, hipStreamCaptureModeGlobal), "begin capture"); + CHECK(vt::rocm::ProductGetBlasStreamIsCapturing(s0) == true); + const hipblasHandle_t h0c = vt::rocm::ProductGetBlasHandle(0, s0); + CHECK(h0c == h0); + CHECK(vt::rocm::ProductGetBlasStreamIsCapturing(s0) == true); + + hipGraph_t graph = nullptr; + RequireHip(hipStreamEndCapture(s0, &graph), "end capture"); + CHECK(vt::rocm::ProductGetBlasStreamIsCapturing(s0) == false); + if (graph != nullptr) { + RequireHip(hipGraphDestroy(graph), "destroy graph"); + } + RequireHip(hipStreamDestroy(s0), "ds0"); +} From aab3b0b32c2b6d5a5a9bba334009e406890ab90f Mon Sep 17 00:00:00 2001 From: Don Mirror Date: Sun, 16 Aug 2026 12:10:58 -0500 Subject: [PATCH 6/7] fix(rocm/gemma4): #837 compile GetBlas product test as HIP Host CXX missed __HIP_PLATFORM_AMD__ on hip_runtime.h. LANGUAGE HIP is the branch contract, not a recipe-only define. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes] --- tests/CMakeLists.txt | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index f880473b2..a236ff421 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1780,6 +1780,8 @@ target_compile_definitions(test_ops_getblas_dualslot PRIVATE VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") if(VLLM_CPP_HIP) # Separate HIP target: calls production GetBlas. Missing devices → exit 77. + # .cpp would compile as host CXX and miss __HIP_PLATFORM_AMD__. + set_source_files_properties(vt/test_ops_getblas_product.cpp PROPERTIES LANGUAGE HIP) vllm_cpp_add_test(test_ops_getblas_product vt/test_ops_getblas_product.cpp) endif() vllm_cpp_add_test(test_ops_penalties vt/test_ops_penalties.cpp) From 090f1c6dc790a4b97725aee52dddd10968b6133a Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Mon, 17 Aug 2026 08:54:57 +0000 Subject: [PATCH 7/7] fix(rocm/gemma4): classify the donor evidence, and leave the FEATURES cell whole Two repairs on top of the #837 dual-slot work. Neither touches its behavior. `.agents/evidence/` is a path `scripts/check-pr-size.py` cannot classify, so the gate refused the change outright rather than sizing it: `ERROR: PR size check could not classify the change: unclassified repository path`. The donor manifest and the extracted `GetBlas` body move to `.agents/specs/` as `rocm-gemma4-getblas-dualslot-donor.md` and `.log`, matching `SPEC` and `SPEC_EVIDENCE` at `check-pr-size.py:188-189`. Teaching the checker a new `.agents/evidence/` pattern would be a governance-checker change owing its own spec and mutation evidence, which this is not. `docs/FEATURES.md` goes back to what main has. The branch had rewritten the Gemma4 row, which dropped the `VT_GEMMA4_*`/`VT_ATTN_*` env pointer, the `test_gemma4_rocm_fp8_seams` seam name and the spec link. Restoring those AND adding the #837 sentence does not fit: `check-public-doc-tables` caps a table cell at 220 characters and main's cell is already 219, which is presumably why the row was shortened in the first place. Since `check-doc-checkpoint` does not ask for a FEATURES edit here -- `src/vt/rocm/` is not a feature surface and no `REGISTER_VLLM_MODEL` set changed (#595, #1086) -- the row keeps its existing content and the #837 detail stays where the branch already put it in full, in `docs/USAGE.md` and the spec. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode] --- .../rocm-gemma4-getblas-dualslot-donor.log} | 0 .../MANIFEST.md => specs/rocm-gemma4-getblas-dualslot-donor.md} | 0 docs/FEATURES.md | 2 +- 3 files changed, 1 insertion(+), 1 deletion(-) rename .agents/{evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt => specs/rocm-gemma4-getblas-dualslot-donor.log} (100%) rename .agents/{evidence/rocm-gemma4-getblas/MANIFEST.md => specs/rocm-gemma4-getblas-dualslot-donor.md} (100%) diff --git a/.agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt b/.agents/specs/rocm-gemma4-getblas-dualslot-donor.log similarity index 100% rename from .agents/evidence/rocm-gemma4-getblas/getblas-fn-67-99.txt rename to .agents/specs/rocm-gemma4-getblas-dualslot-donor.log diff --git a/.agents/evidence/rocm-gemma4-getblas/MANIFEST.md b/.agents/specs/rocm-gemma4-getblas-dualslot-donor.md similarity index 100% rename from .agents/evidence/rocm-gemma4-getblas/MANIFEST.md rename to .agents/specs/rocm-gemma4-getblas-dualslot-donor.md diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 6454c8b47..ebec16432 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -358,7 +358,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | -| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 experts + SharedK-WMMA (RDNA4); #837 GetBlas `tls_slots[2]` + product `GetBlas` probe. Decode-graph deferred | +| Gemma4 MoE ROCm FP8 + SharedK-WMMA | Partial | Dual-GPU FP8 resident experts, SharedK-WMMA prefill (RDNA4); decode-graph and forward extract deferred. Env `VT_GEMMA4_*`/`VT_ATTN_*`, seam `test_gemma4_rocm_fp8_seams`. [spec](../.agents/specs/gemma4-rocm-fp8-moe.md) | ## How to read this page