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..7fcbb8a4c --- /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`. | `IMPLEMENTING` | 2026-08-15 — 6195 production capture hook load-bearing | 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-donor.log b/.agents/specs/rocm-gemma4-getblas-dualslot-donor.log new file mode 100644 index 000000000..6abf5cc0d --- /dev/null +++ b/.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}; +} diff --git a/.agents/specs/rocm-gemma4-getblas-dualslot-donor.md b/.agents/specs/rocm-gemma4-getblas-dualslot-donor.md new file mode 100644 index 000000000..dc69a77a3 --- /dev/null +++ b/.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. diff --git a/.agents/specs/rocm-gemma4-getblas-dualslot.md b/.agents/specs/rocm-gemma4-getblas-dualslot.md new file mode 100644 index 000000000..2526b020f --- /dev/null +++ b/.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`. diff --git a/docs/USAGE.md b/docs/USAGE.md index ab4f6e450..81b6c4385 100644 --- a/docs/USAGE.md +++ b/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 diff --git a/include/vt/rocm/rocm_getblas_dualslot.h b/include/vt/rocm/rocm_getblas_dualslot.h new file mode 100644 index 000000000..8e1b8e70c --- /dev/null +++ b/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 +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 + 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/include/vt/rocm/rocm_getblas_product.h b/include/vt/rocm/rocm_getblas_product.h new file mode 100644 index 000000000..f36ce8aa8 --- /dev/null +++ b/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 +#include + +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 43919747b..a1d043702 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,38 +68,54 @@ 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) { CheckHip(hipSetDevice(d), "hipSetDevice"); } + 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{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; } +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) { return std::string("(") + Name(a.dtype) + "," + Name(b.dtype) + ")->" + Name(out.dtype); } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 4a94639d9..a236ff421 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1775,6 +1775,15 @@ 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}") +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) 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..8c7cf55dd --- /dev/null +++ b/tests/vt/test_ops_getblas_dualslot.cpp @@ -0,0 +1,366 @@ +// #837 GetBlas dual-slot host lifetime + product-call seam (research c24b). +// Production GetBlas and these tests both execute RocmProductGetBlasOn. +#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; + } + } + } +}; + +struct TrackingHooks : FakeHooks { + handle_t Create() { + 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; +}; + +// 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 = 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) { + 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; +} + +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) { + 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") { + 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 = vt::rocm::RocmProductGetBlasOn(eng, 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 via product-call seam") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + TrackingHooks hooks; + Table t; + 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); + 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 = 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); + 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)vt::rocm::RocmProductGetBlasOn(eng, 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 = vt::rocm::RocmProductGetBlasOn(eng, 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; + RunProductHops(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; + RunProductHops(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 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("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 hop 0-1-0-1 restores device on revisit") { + using Eng = vt::rocm::GetBlasDualSlotEngine; + Eng eng; + 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; + 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); + 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..5cdec3386 --- /dev/null +++ b/tests/vt/test_ops_getblas_product.cpp @@ -0,0 +1,114 @@ +// #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"); +} + +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"); +}