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

[bakon11]

Fixes #837.

Problem

On a peer-MoE device hop, the single per-thread (TLS) hipBLAS handle is destroyed
on the wrong device, corrupting the GetBlas accessor during multi-GPU Gemma-4
MoE serving.

What changed

• GetBlas dual-slot TLS keyed by device, plus a host-lifetime seam, so the
  handle is no longer destroyed across the device hop. The production GetBlas
  in src/vt/rocm/rocm_matmul_hipblaslt.hip is edited in place and routes
  through the new engine, so the change is reached rather than merely present.
• Product-call seam and an exit-77 HIP probe; StreamIsCapturing is made
  load-bearing on the product path.
• Spec and identity probe (0→1→0→1) land in the same PR, spec first.

Verification

Built and smoked on gfx1201, shipping ROCm 7.2.4 (clang f58b06d), as the
combined #837+#838+#839 stack on d1b0ea3a: gfx1201 build BUILD_RC=0, all 30
layers resident and bound, ready in ~20 s; batched-MoE prefill deadlock-free
(T=2012 done, PEER_ACT active, no hang, no HIP error); Paris/63 quality PASS.

Stated plainly, because it bounds what this PR alone demonstrates: that smoke
covers the three-PR stack, not this branch in isolation. What is isolated here
is the host-side test evidence — 12 cases with real RED mutants (swapped
selector, destroy-on-hop, missing SetStream, capture-time SetDevice,
forwarded device 0, forwarded null stream), plus a source-invariant case that
reds if the production call site is deleted.

Maintainer changes on top

Two repairs were applied while landing, neither touching behavior:

• The donor evidence moved from .agents/evidence/ — a path
  scripts/check-pr-size.py cannot classify, so the gate refused the change
  outright — to .agents/specs/rocm-gemma4-getblas-dualslot-donor.{md,log},
  which match SPEC and SPEC_EVIDENCE at check-pr-size.py:188-189.
• docs/FEATURES.md is left as main has it. The branch had rewritten the
  Gemma4 row, dropping the VT_GEMMA4_*/VT_ATTN_* env pointer, the
  test_gemma4_rocm_fp8_seams seam name and the spec link. Restoring those and
  adding the ROCm Gemma-4 hyp B: GetBlas single TLS destroys hipBLAS handle on peer-MoE device hop #837 sentence does not fit — check-public-doc-tables caps a cell
  at 220 characters and main's is already 219 — and check-doc-checkpoint does
  not ask for a FEATURES edit here, since src/vt/rocm/ is not a feature
  surface and no REGISTER_VLLM_MODEL set changed (doc-checkpoint makes docs/FEATURES.md a LOCK: 2 of 5 concurrent LTX-2.5 PRs conflict there and nowhere else #595, fix(#595): doc-checkpoint asks whether the registry moved, not whether a file did #1086). The ROCm Gemma-4 hyp B: GetBlas single TLS destroys hipBLAS handle on peer-MoE device hop #837
  detail stays where this branch already put it in full: docs/USAGE.md and the
  spec.

The branch was rebuilt by rebase rather than merge so it carries no untrailered
merge commit; all six original commits are preserved with their authorship.

Known-unrelated CI

windows-msvc-cpu and windows-msvc-vulkan were red on every open PR from a
break predating this branch (#503). That queue has since been cleared — #968
landed as #983 and #1068 as #1069 — and only #584's runtime crash remains.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]