fix(rocm/gemma4): #837 GetBlas dual-slot TLS + host lifetime seam

.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`. | `IMPLEMENTING` | 2026-08-15 — 6195 production capture hook load-bearing |

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

.agents/specs/rocm-gemma4-getblas-dualslot-donor.log

@@ -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};
+}

.agents/specs/rocm-gemma4-getblas-dualslot-donor.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.

.agents/specs/rocm-gemma4-getblas-dualslot.md

@@ -0,0 +1,130 @@
+# 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.
+
+
+## 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`).
+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`.

docs/USAGE.md

@@ -3402,8 +3402,12 @@ 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. `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); 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).
 
 ## LTX-2.5 text conditioning
 

include/vt/rocm/rocm_getblas_dualslot.h

@@ -0,0 +1,75 @@
+// #837 GetBlas dual-slot TLS — slot index + hookable lifetime engine.
+// Product GetBlas and host tests both execute RocmProductGetBlasOn
+// (research c24b: substring wiring is not a load-bearing seam).
+#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 <class Hooks, int (*SlotIndex)(int) = GetBlasSlotIndex>
+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;
+  }
+};
+
+// Product-call seam: production GetBlas and host tests both execute this.
+// Forwards device + stream unchanged. Mutating either argument is RED.
+struct RocmGetBlasForward {
+  template <class Engine, class Hooks>
+  static typename Engine::handle_t apply(Engine& engine, int device,
+                                         typename Engine::stream_t stream,
+                                         Hooks& hooks) {
+    return engine.Get(device, stream, hooks);
+  }
+};
+
+template <class Engine, class Hooks, class Forward = RocmGetBlasForward>
+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

include/vt/rocm/rocm_getblas_product.h

@@ -0,0 +1,18 @@
+// #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 <hip/hip_runtime.h>
+#include <hipblas/hipblas.h>
+
+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