diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 4970ee82a..acd731962 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -23149,3 +23149,108 @@ reads max abs 2.12e-05 / mean rel 2.88e-05, f64-then-store reads 5.88e-08 / below one bf16 ulp of the value it produces. Ten of twelve goldens are unchanged; the two that moved are the V1 arm at 1.19e-07 and 7.45e-08 against a 1e-5 bound. No tolerance was widened. + +## ROCM-GEMMA4-PREFILL-PEER-BARRIER — incremental Finish success-path barrier cost at T=2029 (2026-08-18, #1047 item 3) + +**Attribution only. Not a product ship number. Deleting the two Finish barriers is not authorized to land.** + +Under the frozen dual-R9700 T=2029 prefill-peer recipe, deleting only the two +explicit Finish success-path `hipEventSynchronize(tls.ev_e)` + +`hipStreamSynchronize(cst)` barriers improved end-to-end throughput by 2.55% +(1122.10 vs 1094.24 tok/s; about 46.05 ms/request). This measures the incremental +aggregate cost of those barriers relative to the retained RecordedEvent wait in +`RetirePinThenUnpin`. It does not measure total blocking-retirement cost, prove +asynchronous overlap recovery, or generalize beyond this recipe. + +Researcher verdict `20260818T040236-Researcher-fa20`: active T=2029 timing GREEN. +T=19 remains an optional unused negative control. + +### Construction + +| Arm | Role | Source | BIN sha256 | Size | +|-----|------|--------|------------|------| +| A / BEFORE | wait-only (no extra Finish barriers) | parent `1de7b132` | `4526f3f9122444c8e10d564a7fa87d3ab60b87a5bc6f1cae5e9b411f20a96dca` | 27153608 | +| B / AFTER | two Finish success-path barriers | `cf31e5e98` | `3324394bb1bf1368e86882aa267c2054e53cbbe26dd739f27c636b241b05bf62` | 27153608 | + +Timing BINs are counter-free. Separate identical witness pair (not used for timing): + +| Witness | HEAD | BIN sha256 | Size | +|---------|------|------------|------| +| BEFORE-W | `2b69f6f4` | `a40137930a8652920901abed9be8b71ec65eec8ca4bfef7c9a35b8440dac92a7` | 27155560 | +| AFTER-W | `3aa12542` | `a38e57238920738397140aa90cfd5cabc865ca9bfac2504a27b3b7647020fc71` | 27155560 | + +Harness (frozen at fire): + +| Artifact | SHA256 | +|----------|--------| +| `prefill_peer_client.py` | `adc497419a98e13a1413fea1f64db5060b71900aac945117d7b06061385f027a` | +| `run-1047-item3-timing.sh` | `3c482c6ee6ae27227dd9092cb8d17df916b34fc138c274ac78cc645a2e4ef196` | +| `compare_1047_item3.py` | `e33c2273190f6b3ea5db98e2b6e94b4d06bb972ecb3bd67a9aabbc10b2d38dbc` | +| `parse_prefill_peer_witness.py` | `8ca3cf58c5c192fa142acfad1c24d425d1a3642adb58949de4d3544a40709430` | + +### Recipe + +Isolated `:8012` only. Never `:8010`. Dual Radeon AI PRO R9700 (gfx1201). Model +`gemma-4-26B-A4B-it-fp8`. Order calib then A1 -> B1 -> B2 -> A2, separate +processes, teardown between each. + +Env: `HIP_VISIBLE_DEVICES=0,1` `PREFIX_CACHE=0` `MAX_MODEL_LEN=65536` +`NUM_BLOCKS=2048` `MAX_NUM_SEQS=4` `MAX_BATCHED_TOKENS=8192` +`VT_GEMMA4_PREFILL_PEER_ACT=1` `VT_GEMMA4_SWA_PHYSICAL=0` +`VT_GEMMA4_GPU0_HEADROOM_GB=16`. `VT_GEMMA4_PREFILL_PEER_WITNESS` and +`VT_GEMMA4_PROFILE` unset on every timing arm. + +Counter-free calibration on the BEFORE timing BIN froze prompt SHA +`bdd0bfac665d250ede60f1f49f6a77106736afac01a80c559bf2143142ec54b8` +(`usage.prompt_tokens=2029` exact). Same file reused on all four legs. + +Identity gate (all four legs): HTTP 200, `prompt_tokens=2029`, +`completion_tokens=2`, tokens `["**","Count"]`, `seq_sha=cf33d9bd3b418e54e49f2358`. + +Each timing leg: 3 discarded warmups, 5 measured bursts (`burst=1`, +`out-tokens=1`). `n_tok` from API `usage` only. `dirty_excluded=0`, `resets=[]`. +Zero `prefill_peer_witness:` lines in timing serve logs. + +### 20 raw measured samples (tok_sum=2030 each) + +| leg | i0 | i1 | i2 | i3 | i4 | +|-----|---:|---:|---:|---:|---:| +| A1 | 1125.2743 | 1122.6271 | 1122.5586 | 1121.6334 | 1119.2061 | +| B1 | 1092.3775 | 1095.8747 | 1092.4946 | 1092.6152 | 1089.7476 | +| B2 | 1096.5926 | 1099.3896 | 1084.0359 | 1096.2230 | 1098.4552 | +| A2 | 1120.4249 | 1117.9151 | 1123.3178 | 1121.0147 | 1125.8531 | + +### Comparator (independent recompute matches) + +- A pool median **1122.0960** tok/s, range 7.9380 +- B pool median **1094.2450** tok/s, range 15.3537 +- A/B **1.025452x** (+2.5452% without the barriers) +- Median request wall: A **1.809115 s**, B **1.855165 s**; delta **46.050 ms** +- Effect **27.8510** tok/s; max pooled arm range **15.3537**; effect/max-range **1.814x**; `inside_disp=false` +- Leg medians keep direction across the interleave: A1 1122.5586, B1 1092.4946, B2 1096.5926, A2 1121.0147 +- A1<->A2 drift 1.5439; B1<->B2 drift 4.0980 +- dirty A1=0 B1=0 B2=0 A2=0 +- `PROTOCOL_OK`; comparator printed `CANDIDATE_PROMOTE`; lab did not self-promote + +### Live witness (equal-call proof, unmeasured) + +Event-derived each arm: `n_lines=4844`, `launch_ok=finish_ok=2422`, fails=0, +`y_host_fb=0`. A/B M histograms identical (639 keys, sum 2422, M in [1,2020]). +`witness-compare.json` sha256 +`74ae4180db5a3d0b499c371932f71e4fa50f2cc31b267fafb394f0a689a38101`. +WITNESS absent during calib, present on both witness arms. + +### Artifacts + +`/home/don/.cache/hermes-builds/pr1047-cost/ab-out-item3-timing/` +(orchestrator.log, compare.txt, timing-arm{A1,B1,B2,A2}-T2029.json, +ident-arm*-T2029.json, serve-arm*.log, environ-arm*.txt, frozen prompt). +Witness pack: `ab-out-item3-witness/`. Bus RESULT `20260818T040119-hermes-1f25`. + +### What this is NOT + +- Not a license to delete the Finish waits in product. +- Not total blocking-retirement cost vs a no-host-wait baseline. +- Not overlap / async-unpin recovery. +- Not a KEEP `:8010` or p42k number. +- Not T=19 (control not authorized). diff --git a/.agents/claims/CLAIM-ROCM-GEMMA4-PREFILL-PEER-HELPER.md b/.agents/claims/CLAIM-ROCM-GEMMA4-PREFILL-PEER-HELPER.md new file mode 100644 index 000000000..b5555bf75 --- /dev/null +++ b/.agents/claims/CLAIM-ROCM-GEMMA4-PREFILL-PEER-HELPER.md @@ -0,0 +1,5 @@ +# CLAIM-ROCM-GEMMA4-PREFILL-PEER-HELPER + +| Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-ROCM-GEMMA4-PREFILL-PEER-HELPER` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-PREFILL-PEER-HELPER`, issue #839) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-prefill-peer` | `row/ROCM-GEMMA4-PREFILL-PEER-HELPER` | Owns ONLY: `RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice` → Launch/Finish + PeerPipeTls + DequantCache pin that persists in PeerSlot until host-observed retirement (ev_e or expert-stream sync on pre-record errors); Ensure rejects reconfiguration while pins live. **EXCLUDED:** kPeerPipe default ON, FP8 Lt, GU_INTERLEAVE, donor unpin-before-ev_e, donor Ensure→FreeAll on live pins, #838, #697. Requires #837 before or with impl as a separate head. Independent history from abandoned `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — ea9c: product RestoreComputeOrThrow no-op mutation RED via HIP compile-and-run | diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 6a8340284..dc287903c 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -365,6 +365,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1171](https://github.com/mudler/vllm.cpp/issues/1171) | `KERNEL-GDN-REPLAYSSM` | GDN decode rewrites the whole `[HV,V,K]` fp32 state every step (`src/vt/cuda/cuda_gdn.cu:2393` reads the tile, `:2425` writes it back), which at the 27B shape `HV=32, V=128, K=128` is 2,097,152 bytes read plus the same written per layer, per request, per token. ReplaySSM keeps a per-slot ring of the last `L` steps' rank-1 factors `(d, k, g)`, reconstructs the state in registers, and writes it back only every `L` steps. vLLM implements the algorithm at the pin `555967922` for Mamba2 selective-state ONLY (`layers/mamba/ops/selective_state_update_replayssm_output_only.py`, ring shapes/dtypes `mamba_utils.py:84-93,202-221`, `use_replayssm` default `False` at `config/cache.py:152`, introduced `866fea2b` #48018) and it cannot reach GDN: `config/vllm.py:2318-2322` refuses any model not setting `supports_replayssm`, only `NemotronHForCausalLM` sets it (`models/nemotron_h.py:711`), `GDNAttentionMetadataBuilder` does not subclass the Mamba builder that derives the cursor (`v1/attention/backends/gdn_attn.py:82` vs `mamba_attn.py:575-638`), and the kernel hard-requires a scalar-per-head `A` (`:540-542`) with the Mamba2 `(B,C)` group structure (`:529`). Still true 877 commits past the pin. SGLang ported it to GDN at our recorded pin `f63458b5` (`layers/attention/fla/fused_recurrent_linear_replayssm.py`, whose `:50` credits vLLM; `--enable-linear-replayssm` default `False` and `--linear-replayssm-cache-len` default 16 at `server_args.py:1972-1986`; rings `memory_pool.py:465-483`; commit `a10a24e9` #28451), so the algorithm is a vLLM mirror and the GDN application is a secondary-oracle port. PAYOFF UNMEASURED HERE and deliberately not claimed: ReplaySSM removes the state WRITE and not the read, the flush step reads the checkpoint a SECOND time, so the honest state ratio is `(1+2/L)/2` = 0.5625 at `L=16` against SGLang's published 0.53x which models neither the flush re-read nor any ring read; the ring itself adds 395,264 bytes per slot per layer at the 27B shape = **+18.9% KV page**, worse than vLLM's ~7% on Nemotron because GDN's state is `V*K` while the ring is `L*(V+K)`; and SGLang's own end-to-end figure is ~2.3% TPOT at 128 concurrency on an MoE model. Neither upstream is bit-exact against its unbuffered path and neither claims to be. Motivation is the open Qwen3.8-27B bf16 decode gap (c4 total 0.918x, output 0.963x, `docs/BENCHMARKS.md:192-205`). Spec [`gdn-replayssm.md`](specs/gdn-replayssm.md) | perf | | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | `ENG-CUDAGRAPH-BREAK` | The hand-rolled decode-graph driver count recorded in `9bc4d7f44` is **eight** and is actually **nine**, and the row it feeds was framed as coverage-only when it is also correctness. The ninth is the DFlash draft graph, file-local with no header declaration, at `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106` — its own `int g_state = 0` three-state machine (`:771`), its own `VT_DFLASH_GRAPH` kill switch (`:870`) instead of the `VLLM_CPP_CUDAGRAPH` the six batched drivers read, its own invalidate-on-block-width-change (`:1038-1047`) and its own `try { EndCaptureGraph(); } catch (...) {}` drain (`:1106`). The eight-count is stated in four places, all corrected here: [`sglang-breakable-cuda-graph.md`](specs/sglang-breakable-cuda-graph.md) §4 and `## Owed`, [`.agents/engine-matrix.md`](engine-matrix.md) rows `ENG-CUDAGRAPH-BREAK` and `ENG-CUDAGRAPH-DEDUP` ("times eight drivers", which sizes #1162's signature table), and [`.agents/roadmap_v1.md`](roadmap_v1.md) track `C12`. The reframing is the substantive half: `ENG-CUDAGRAPH-BREAK` was recorded as a COVERAGE row, and the duplication has already cost a SHIPPED model its decode graph. `src/vllm/model_executor/models/qwen3.cpp:961-986` declines the decode graph outright whenever the asynchronous device-token mirror is live, on its own measured battery — `depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate` — because `Step()` replays against the HOST `input.token_ids` and the combine has patched the DEVICE ids. The comment names the real fix as reading the identifiers at replay time from a stable device buffer, and that fix exists, in exactly one sibling driver, as `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`): `grep -c StepDevInputs` returns 41 lines there and 0 in each of `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`. One capability, written once, unavailable to four models, with a live mitigation standing in its place. This does NOT weaken the framing rule that `ENG-CUDAGRAPH` established: the row still makes no throughput claim, and the prefill refutation (GB10 3.8% host-idle between launches, GPU-busy >96%, 27B prefill gap 92.5% non-GEMM glue) stands unchanged. Coverage AND correctness, never speed. Fixed in flow with the [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) review repair ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)) | record | | [#1181](https://github.com/mudler/vllm.cpp/issues/1181) | `FIX-READ-F32-SCALAR-GUARD` | `ReadF32Scalar` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:312-318` @ `ab6e65216`) bounds its input with `t.data != nullptr && t.nbytes >= sizeof(float)`, a LOWER bound, and then `memcpy`s four bytes into a `float`. Two silent wrong-value paths follow and neither fails: an ARRAY is reduced to element 0, so a block-wise FP8 scale grid of shape `[ceil(N/128), ceil(K/128)]` passes and stands in for the whole weight (measured under [#1166](https://github.com/mudler/vllm.cpp/issues/1166) on `Qwen/Qwen3.8-27B-FP8` @ `017b9c7af6b5689d5dd426a76e0bc077eb5ca20a`, `q_proj.weight_scale_inv` is `[96, 40]`), and ANY dtype is reinterpreted, since that same tensor is `BF16` and its four bytes are two bf16 values read as one float. Both return a finite plausible float, so the output is fluent, plausible and wrong, which is what a token gate cannot see. Upstream makes both facts structural rather than optional: a per-tensor scale is a distinct parameter TYPE that asserts `loaded_weight.shape[0] == 1` (`vllm/model_executor/parameter.py:260-272,304-309` @ `555967922`, plus the `_assert_and_load` shape assert at `:93-96`), the slot is allocated `torch.float32` so a narrow on-disk dtype is VALUE-converted rather than reinterpreted (`utils/fp8_utils.py:1276`), and the declared strategy TENSOR/CHANNEL/BLOCK picks the parameter type before a byte is read (`compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py:63,128`). The AUDIT corrects the issue's own framing twice. The 27 grep hits across five files are 5 definitions, 20 call sites and 2 comment references, and both counts are short: `ReadCtF32Scalar` (`include/vllm/model_executor/models/dense_weight_loaders.h:376`) is a SIXTH copy of the same defect under another name, reached from a SIXTH model file (`src/vllm/model_executor/models/qwen3_weights.cpp:100,126-128` through `LoadCtNvfp4W4A16`). Of the six, three check nothing, `LnReadF32Scalar`/`ShReadF32Scalar` check dtype but not count, and only `nemotron_h_weights.cpp:557-573` is correct, which makes it the model the shared guard generalizes. No call site legitimately passes a multi-element or non-F32 tensor, and every existing fixture emits rank-0 or `{1}` `F32`, so nothing in the tree needed the leniency. It is NOT merely latent: `dense_weight_loaders.h:73-74` and `docs/BENCHMARKS.md:52` both record `unsloth/Qwen3.6-27B-NVFP4` @ `ccdaab7e` as FP8 W8A8 throughout with BF16 PER-OUTPUT-CHANNEL scales, and `LoadAttnDense` branches on the weight dtype alone (`qwen3_5_dense_weights.cpp:478-480`), so those projections enter the per-tensor arm and hit both defects at once under the tensor name the loader actually asked for, with no misspelling to stop them. Fixed in flow by one `dense_loaders::ReadF32Scalar(get, name)` that refuses `numel != 1` naming the shape, refuses a non-`F32` dtype naming the dtype, and requires exactly four readable bytes, with the other five copies deleted onto it and `nemotron_h`'s `Loader`-based twin kept as the one tracked exception. A narrow dtype is refused rather than converted, because a one-element BF16 scale has never been read correctly here and the BF16 layout that IS shipped is per-channel, which the count check refuses first. Per-channel FP8, block-wise FP8 and any explicit narrow-dtype conversion stay owed. Spec [`read-f32-scalar-guard.md`](specs/read-f32-scalar-guard.md) | bug | +| [#839](https://github.com/mudler/vllm.cpp/issues/839) | `ROCM-GEMMA4-PREFILL-PEER-HELPER` | Gemma4 prefill peer helper: Launch/Finish pin until host-observed retire | perf | | [#1187](https://github.com/mudler/vllm.cpp/issues/1187) | — | The image-conditioning request surface is TWO FIXED SLOTS against upstream's repeatable `--image PATH FRAME_IDX STRENGTH [CRF]` (`packages/ltx-pipelines/src/ltx_pipelines/utils/args.py:805-817` @ `fd4ded7f`), which both conditioning builders expand per item (`utils/helpers.py:272-308` and `:343-367`). This ABI carries `first_frame` / `first_frame_ppm` and `last_frame` (`include/vllm.h:947-948`) at the FIXED indices `latent_idx = 0` and `frame_idx = frames - 1`, with one shared `noise_aug` and one shared `image_crf`, so an interior keyframe, a third keyframe, and two keyframes at different strengths or CRFs are each unrequestable; `ltx2-gen` exposes only `--first-frame` and `/v1/videos` forwards only a first frame. `num_generated_keyframes` is a DIFFERENT feature — model-invented interior slots (`keyframe_slots.py:71-150`) — and is not a substitute. SECOND HALF, filed together because it is the same surface: the sigma-BINNED guider factory. `FactoryGuidedDenoiser` resolves a guider per step from sigma (`utils/denoisers.py:332-343`) through `MultiModalGuiderFactory.build_from_sigma` (`ltx-core/components/guiders.py:337-342`), but every `main()` in `ltx-pipelines` passes plain `MultiModalGuiderParams`, so `create_multimodal_guider_factory` takes `constant()` and builds ONE `(inf, params)` bin (`guiders.py:312-315`) — which is exactly what `Ltx2GuidedDenoise` mirrors. The unported arm is `MultiModalGuiderFactory.from_dict` (`guiders.py:317-330`), reachable only by a caller who builds a factory and passes it as `video_guider_params`; this engine holds one `Ltx2MultiModalGuiderParams` per phase and has no sigma-keyed spelling. This CORRECTS the framing of [#1096](https://github.com/mudler/vllm.cpp/issues/1096), which named the per-sigma denoiser as a blocker on `KeyframeInterpolationPipeline` without deriving that the pipeline's own default path never resolves more than one guider. Filed rather than fixed in flow because it changes `include/vllm.h`, the video-engine request struct, `ltx2-gen`, the server request mapping and the conditioning loop, needs its own goldens for per-item strength and CRF, and narrows EVERY pipeline that takes `images` rather than one recipe. Listed under `## Owed` in [`ltx25-keyframe-interp.md`](specs/ltx25-keyframe-interp.md) | feature | | [#1191](https://github.com/mudler/vllm.cpp/issues/1191) | `LTX25-KEYFRAME-INTERP` | `ltx2-gen` parsed `--first-frame` (`examples/ltx2_gen/main.cpp:333`) and assigned `vp.first_frame` (`:413`) while never reading `last_frame` at all, though `vllm_video_params` has carried the field since `include/vllm.h:948` and the LTX-2.5 engine has SERVED it since row `LTX25-TOKEN-APPEND` ([#930](https://github.com/mudler/vllm.cpp/issues/930)): the `wants_last_frame` arm of the phase loop places the image as a `VideoConditionByKeyframeIndex` at pixel frame `frames - 1`, gated by `test_ltx2_video`. So a closing keyframe was reachable from the C API and not from the shipped CLI. Latent until row `LTX25-KEYFRAME-INTERP` ([#1096](https://github.com/mudler/vllm.cpp/issues/1096)), which is the first caller it actually narrows: `KeyframeInterpolationPipeline` exists to generate the motion BETWEEN pinned keyframes, and with one slot the CLI can only ask for half of that. Fixed in flow by that row — one `--last-frame` flag parsed and assigned beside `--first-frame`, sharing the same `--image-crf` and strength the two slots already share, plus the usage and help text. An INTERIOR keyframe stays unrequestable and is [#1187](https://github.com/mudler/vllm.cpp/issues/1187) | bug | | [#1185](https://github.com/mudler/vllm.cpp/issues/1185) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle BUILDS, installs, imports and sees the GPU inside an `rc` lease on `dgx:gpu0`, measured 2026-08-18, which falsifies the `nvcc` clause four records carried. [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) said the oracle "needs `nvcc`, which the worker still lacks", and `.agents/environment.md`, [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) and [`gpu-lease-methodology.md`](specs/gpu-lease-methodology.md) each derived a blocker from it. The build job (`buildvllm.sh`, staged sha256 prefix `15e140d41f44e7c2`) asserted the checkout against the pin BEFORE compiling, printing `PIN CONFIRMED` at `5559679229bc961848b121ccdeaa8fa5d79bec98` and aborting otherwise, took `nvcc` from the toolkit row `MODEL-NEMOTRON-H-ABI-A3-E2E` staged (`NVCC_RC=0`, CUDA `release 13.3, V13.3.73`) and produced `WHEEL_RC=0`, `PERSIST_RC=0` and a 434 MiB `vllm-0.1.dev1+g555967922.cu133-cp312-cp312-linux_aarch64.whl`, sha256 `7c58b339741a288fbb313f4f5196c9c92a9e3b3c3ebe2ea970b0ff50bb9bcba4`. The identity job (`oracleenv.sh`, prefix `6119f5223f5d818c`) asserted from `cd /`, outside any source tree: `vllm.__version__ = 0.1.dev1+g555967922`, `IDENTITY_RC=0`, `cuda True NVIDIA GB10`, `CUDA_RC=0`. SCOPE, and it carries the same weight as the result: RUNNING A MODEL IS UNTESTED. Only build, install, import and `torch.cuda.is_available()` are measured, and [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) records that the last time an oracle reached this far it consumed the host in the step AFTER `torch.compile` and REBOOTED the box, at `gpu_memory_utilization` 0.75 and again at 0.30, so the fraction is not the lever. The version string is an OPEN discrepancy: `.agents/upstream-sync.md` records `vllm_runtime_version = 0.23.1rc1.dev1511+g555967922`, the commit segment matches and satisfies the pin's binding `+g` rule, and the prefix differs because a shallow fetch stops `setuptools_scm` counting commits since the last tag, so a full-string gate needs a deeper fetch or a recorded pretend-version. The venv is NOT staged, because that job was killed at a 90-minute ceiling mid-copy and its partial tree was removed, so only the WHEEL is durable. Four staging walls, all artifacts of the NAS rather than of CUDA: `cp -a` preserves `file_mode=0664` so `nvcc` exited 126. CIFS `nounix` stores no symlink so `include` and `lib64` vanished and CMake reported `Could NOT find CUDA (missing: CUDA_INCLUDE_DIRS CUDA_CUDART_LIBRARY) (found version "13.3")`, naming the version and denying the toolkit in one line. 32 library links `libfoo.so` and `libfoo.so.MAJOR` had to be rebuilt because only the `libfoo.so.X.Y.Z` real files survived. And `markupsafe` existed as a dist-info with NO package files from a `pip --target` killed at a 35-minute ceiling, so Marlin codegen died on `ModuleNotFoundError`. The `rc` worker container is REUSED between jobs, so a repair inside a staging branch is skipped on the next run (`nvcc already in place`) and an environment repair must be unconditional and assert its postcondition. CONSEQUENCE for the rows #1129 blocked, [#1003](https://github.com/mudler/vllm.cpp/issues/1003), [#915](https://github.com/mudler/vllm.cpp/issues/915), [#821](https://github.com/mudler/vllm.cpp/issues/821) and [#81](https://github.com/mudler/vllm.cpp/issues/81): UNBLOCKED FOR THE BUILD STEP and STILL BLOCKED FOR A MODEL RUN. None can take a measurement until a model run is demonstrated. Job details, walls and non-claims in [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) | verification | diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-dequant-cache-210-232.log b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-dequant-cache-210-232.log new file mode 100644 index 000000000..65b43f385 --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-dequant-cache-210-232.log @@ -0,0 +1,23 @@ + } + + void UnpinLocked(int idx) { + if (idx < 0 || idx >= nslots) return; + if (slots[idx].pins > 0) slots[idx].pins--; + } + + std::mutex mu; +}; + +// Process-wide per-device caches (NOT thread_local — engine worker threads would +// each thrash a private 1-slot-effective working set and never hit). +struct DequantCacheSlot { + PrefillDequantCache cache; +}; +DequantCacheSlot& DequantCacheSlotFor(int device) { + static DequantCacheSlot slots[8]; + if (device < 0 || device >= 8) { + static DequantCacheSlot fallback; + return fallback; + } + return slots[device]; +} diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-dequant-cache-core-60-209.log b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-dequant-cache-core-60-209.log new file mode 100644 index 000000000..3cb99eedf --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-dequant-cache-core-60-209.log @@ -0,0 +1,150 @@ +// Multi-expert sticky FP8→BF16 weight cache (OPT-IN). +// rocprof p11k: Fp8ChannelDequant ~10–18% device time. Single-slot thrash is real +// within a layer's multi-chunk GEMM, BUT each MoE layer owns distinct expert weight +// tensors — uniq_keys grows to ~O(layers×experts) (~1k+ @ p11k). A small LRU cannot +// absorb cross-layer dequant without multi-GB BF16 twins. Default OFF (1 slot). +// VT_GEMMA4_PREFILL_DEQUANT_CACHE = slot count (default 1, max 128). +inline int PrefillDequantCacheSlots() { + static const int n = [] { + if (const char* e = std::getenv("VT_GEMMA4_PREFILL_DEQUANT_CACHE")) { + int v = std::atoi(e); + if (v <= 0) return 1; + return std::min(v, 128); + } + return 1; // legacy single-slot; multi-slot is lab opt-in only + }(); + return n; +} + +// Per-device LRU of dequantized [gu_w, dn_w] keyed by stable fp8_gu (or weight_id). +struct PrefillDequantCache { + struct Slot { + const void* key = nullptr; + void* gu = nullptr; + void* dn = nullptr; + uint64_t lru = 0; + int pins = 0; + }; + int dev = -1; + int I = 0, H = 0; + int nslots = 0; + uint64_t clock = 0; + uint64_t hits = 0, misses = 0; + Slot slots[128]{}; + + void FreeAll() { + if (dev >= 0) (void)hipSetDevice(dev); + for (int i = 0; i < nslots; ++i) { + if (slots[i].gu) { + (void)hipFree(slots[i].gu); + slots[i].gu = nullptr; + } + if (slots[i].dn) { + (void)hipFree(slots[i].dn); + slots[i].dn = nullptr; + } + slots[i].key = nullptr; + slots[i].lru = 0; + slots[i].pins = 0; + } + nslots = 0; + dev = -1; + I = H = 0; + } + + bool Ensure(int device, int i_dim, int h_dim) { + const int want = PrefillDequantCacheSlots(); + if (dev == device && I == i_dim && H == h_dim && nslots == want) return true; + FreeAll(); + if (device < 0 || i_dim <= 0 || h_dim <= 0 || want <= 0) return false; + if (hipSetDevice(device) != hipSuccess) return false; + const size_t gu_b = static_cast(2 * i_dim) * static_cast(h_dim) * 2; + const size_t dn_b = static_cast(h_dim) * static_cast(i_dim) * 2; + for (int i = 0; i < want; ++i) { + if (hipMalloc(&slots[i].gu, gu_b) != hipSuccess) { + FreeAll(); + return false; + } + if (hipMalloc(&slots[i].dn, dn_b) != hipSuccess) { + FreeAll(); + return false; + } + } + dev = device; + I = i_dim; + H = h_dim; + nslots = want; + clock = 0; + return true; + } + + // On hit: returns cached BF16 weights. On miss: dequant into LRU victim. + // Caller must hold mu. Sets *pin_out so caller can Unpin after GEMM. + bool GetLocked(vt::Queue& q, const void* key, const void* fp8_gu, const void* s_gu, + const void* fp8_dn, const void* s_dn, void** gu_out, void** dn_out, + int* pin_out) { + if (!key || !fp8_gu || !gu_out || !dn_out || nslots <= 0) return false; + int hit = -1; + int victim = -1; + uint64_t oldest = UINT64_MAX; + for (int i = 0; i < nslots; ++i) { + if (slots[i].key == key) { + hit = i; + break; + } + // Never evict a pinned slot. + if (slots[i].pins > 0) continue; + if (slots[i].key == nullptr) { + if (victim < 0) victim = i; + } else if (slots[i].lru < oldest) { + oldest = slots[i].lru; + if (victim < 0 || slots[victim].key != nullptr) victim = i; + } + } + if (hit >= 0) { + ++hits; + slots[hit].lru = ++clock; + slots[hit].pins++; + *gu_out = slots[hit].gu; + *dn_out = slots[hit].dn; + if (pin_out) *pin_out = hit; + return true; + } + if (victim < 0) { + // All pinned — fall back to any unpinned or slot 0 + for (int i = 0; i < nslots; ++i) { + if (slots[i].pins == 0) { + victim = i; + break; + } + } + if (victim < 0) return false; + } + ++misses; + Slot& s = slots[victim]; + vt::DequantFp8ChannelBf16(q, s.gu, fp8_gu, s_gu, 2 * I, H); + vt::DequantFp8ChannelBf16(q, s.dn, fp8_dn, s_dn, H, I); + s.key = key; + s.lru = ++clock; + s.pins++; + *gu_out = s.gu; + *dn_out = s.dn; + if (pin_out) *pin_out = victim; + // One-time unique-key census (lab) + { + static std::mutex uk_mu; + static std::unordered_set uniq; + static int uk_logs = 0; + std::lock_guard ulk(uk_mu); + uniq.insert(key); + if (uk_logs < 3 || misses == 64 || misses == 128 || misses == 256 || + (misses % 2048ull) == 0ull) { + std::fprintf(stderr, + "INFO gemma4-moe dequant_cache dev=%d slots=%d miss=%llu hit=%llu " + "uniq_keys=%zu key=%p\n", + dev, nslots, static_cast(misses), + static_cast(hits), uniq.size(), key); + ++uk_logs; + } + } + return true; diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-launch-finish-1090-1310.log b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-launch-finish-1090-1310.log new file mode 100644 index 000000000..c0c89aa44 --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-launch-finish-1090-1310.log @@ -0,0 +1,221 @@ + + // PEER path: dual-slot Launch+Finish (slot 0). Pipeline uses Launch/Finish directly. + if (!LaunchGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0, expert_dev, x_compute, fp8_gu, + s_gu, fp8_dn, s_dn, M, I, H)) + return false; + return FinishGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0, y_compute, M, H); +} + +bool LaunchGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, int expert_dev, + const void* x_compute, const void* fp8_gu, + const void* s_gu, const void* fp8_dn, const void* s_dn, + int M, int I, int H) { + if (slot < 0 || slot > 1 || !x_compute || !fp8_gu || !s_gu || !fp8_dn || !s_dn || M <= 0 || + I <= 0 || H <= 0 || expert_dev < 0) + return false; + const int compute_dev = compute_q.device.index; + if (compute_dev < 0 || expert_dev == compute_dev) return false; + + static const bool prefill_fp8_lt = [] { + if (const char* e = std::getenv("VT_GEMMA4_PREFILL_FP8_LT")) return e[0] == '1'; + return false; + }(); + static const bool gu_pair = [] { + const char* e = std::getenv("VT_GEMMA4_GU_INTERLEAVE"); + return e && e[0] == '1'; + }(); + + PeerSlot& tls = PeerPipeTls().s[slot]; + + auto freep = [](int dev, void*& p) { + if (!p) return; + (void)hipSetDevice(dev); + (void)hipFree(p); + p = nullptr; + }; + const size_t row_b = static_cast(H) * 2; + const size_t xy_need = static_cast(M) * row_b; + vt::Device edev{vt::DeviceType::kROCM, expert_dev}; + hipStream_t cst = static_cast(compute_q.handle); + + if (tls.pending_M > 0) { + std::fprintf(stderr, "WARN gemma4-moe PEER Launch slot=%d still pending\n", slot); + return false; + } + + if (tls.edev != expert_dev || tls.cdev != compute_dev || tls.I != I || tls.H != H || + tls.Mcap < M || !tls.x) { + freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.x); + freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.y); + freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.gu_act); + freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.act); + if (tls.eq_live) { + vt::DestroyQueue(tls.eq); + tls.eq_live = false; + } + if (tls.ev_c) { + (void)hipEventDestroy(tls.ev_c); + tls.ev_c = nullptr; + } + if (tls.ev_e) { + (void)hipEventDestroy(tls.ev_e); + tls.ev_e = nullptr; + } + const int Mcap = std::max(M, 64); + Check(hipSetDevice(expert_dev), "peer-pipe expert alloc"); + if (hipMalloc(&tls.x, static_cast(Mcap) * row_b) != hipSuccess) return false; + if (hipMalloc(&tls.y, static_cast(Mcap) * row_b) != hipSuccess) return false; + if (hipMalloc(&tls.gu_act, static_cast(Mcap) * static_cast(2 * I) * 2) != + hipSuccess) + return false; + if (hipMalloc(&tls.act, static_cast(Mcap) * static_cast(I) * 2) != hipSuccess) + return false; + tls.eq = vt::CreateQueue(edev); + tls.eq_live = true; + Check(hipSetDevice(compute_dev), "peer-pipe ev_c"); + Check(hipEventCreateWithFlags(&tls.ev_c, hipEventDisableTiming), "ev_c"); + Check(hipSetDevice(expert_dev), "peer-pipe ev_e"); + Check(hipEventCreateWithFlags(&tls.ev_e, hipEventDisableTiming), "ev_e"); + Check(hipSetDevice(compute_dev), "p2p c"); + EnablePeerAccessQuiet(expert_dev); + Check(hipSetDevice(expert_dev), "p2p e"); + EnablePeerAccessQuiet(compute_dev); + tls.edev = expert_dev; + tls.cdev = compute_dev; + tls.I = I; + tls.H = H; + tls.Mcap = Mcap; + } + + hipStream_t est = static_cast(tls.eq.handle); + if (hipEventRecord(tls.ev_c, cst) != hipSuccess) return false; + if (hipStreamWaitEvent(est, tls.ev_c, 0) != hipSuccess) return false; + + bool x_ok = hipMemcpyPeerAsync(tls.x, expert_dev, const_cast(x_compute), compute_dev, + xy_need, est) == hipSuccess; + if (!x_ok) { + std::vector tmp(xy_need); + Check(hipSetDevice(compute_dev), "x d2h"); + if (hipStreamSynchronize(cst) != hipSuccess) return false; + if (hipMemcpy(tmp.data(), x_compute, xy_need, hipMemcpyDeviceToHost) != hipSuccess) return false; + Check(hipSetDevice(expert_dev), "x h2d"); + if (hipMemcpyAsync(tls.x, tmp.data(), xy_need, hipMemcpyHostToDevice, est) != hipSuccess) + return false; + } + + Check(hipSetDevice(expert_dev), "geglu setdev"); + auto try_fp8_lt = [&](vt::Queue& q, void* y, const void* x, int m) -> bool { + if (!prefill_fp8_lt || m < 16) return false; + struct LtTls { + int dev = -1; + int I = 0, H = 0, Mcap = 0; + void* gu_act = nullptr; + void* act = nullptr; + }; + static thread_local LtTls lt; + const int dev = q.device.index; + if (lt.dev != dev || lt.I != I || lt.H != H || lt.Mcap < m || !lt.gu_act) { + auto freep2 = [&](void*& p) { + if (p) { + (void)hipFree(p); + p = nullptr; + } + }; + freep2(lt.gu_act); + freep2(lt.act); + Check(hipSetDevice(dev), "lt alloc"); + const int Mcap = std::max(m, 64); + if (hipMalloc(<.gu_act, static_cast(Mcap) * static_cast(2 * I) * 2) != + hipSuccess) + return false; + if (hipMalloc(<.act, static_cast(Mcap) * static_cast(I) * 2) != hipSuccess) + return false; + lt.dev = dev; + lt.I = I; + lt.H = H; + lt.Mcap = Mcap; + } + if (!vt::MatmulBTFp8ActWeightLt(q, lt.gu_act, x, fp8_gu, s_gu, m, 2 * I, H, 1.f, 0.f)) + return false; + vt::Tensor act = vt::Tensor::Contiguous(lt.act, vt::DType::kBF16, q.device, {m, I}); + vt::Tensor gu_act = vt::Tensor::Contiguous(lt.gu_act, vt::DType::kBF16, q.device, {m, 2 * I}); + if (gu_pair) + vt::rocm::GeluAndMulPairRocm(q, act, gu_act); + else + vt::GeluAndMul(q, act, gu_act); + if (!vt::MatmulBTFp8ActWeightLt(q, y, lt.act, fp8_dn, s_dn, m, H, I, 1.f, 0.f)) return false; + return true; + }; + + bool lt_ok = try_fp8_lt(tls.eq, tls.y, tls.x, M); + if (!lt_ok) { + auto& wslot = DequantCacheSlotFor(expert_dev); + void* gu_w = nullptr; + void* dn_w = nullptr; + int pin = -1; + { + std::lock_guard lk(wslot.cache.mu); + if (!wslot.cache.Ensure(expert_dev, I, H)) return false; + Check(hipSetDevice(expert_dev), "dequant setdev"); + if (!wslot.cache.GetLocked(tls.eq, fp8_gu, fp8_gu, s_gu, fp8_dn, s_dn, &gu_w, &dn_w, &pin)) + return false; + } + vt::Tensor x = vt::Tensor::Contiguous(tls.x, vt::DType::kBF16, edev, {M, H}); + vt::Tensor gu_wt = vt::Tensor::Contiguous(gu_w, vt::DType::kBF16, edev, {2 * I, H}); + vt::Tensor gu_act = vt::Tensor::Contiguous(tls.gu_act, vt::DType::kBF16, edev, {M, 2 * I}); + vt::Tensor act = vt::Tensor::Contiguous(tls.act, vt::DType::kBF16, edev, {M, I}); + vt::MatmulBT(tls.eq, gu_act, x, gu_wt); + if (gu_pair) + vt::rocm::GeluAndMulPairRocm(tls.eq, act, gu_act); + else + vt::GeluAndMul(tls.eq, act, gu_act); + vt::MatmulBTAlphaBeta(tls.eq, tls.y, act.data, dn_w, M, H, I, 1.f, 0.f, vt::DType::kBF16); + { + std::lock_guard lk(wslot.cache.mu); + wslot.cache.UnpinLocked(pin); + } + } + + if (hipEventRecord(tls.ev_e, est) != hipSuccess) return false; + tls.pending_M = M; + if (hipSetDevice(compute_dev) != hipSuccess) return false; + return true; +} + +bool FinishGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, void* y_compute, int M, + int H) { + if (slot < 0 || slot > 1 || !y_compute || M <= 0 || H <= 0) return false; + PeerSlot& tls = PeerPipeTls().s[slot]; + if (tls.pending_M <= 0 || !tls.y || !tls.ev_e) return false; + if (M > tls.pending_M) return false; + const int compute_dev = compute_q.device.index; + const int expert_dev = tls.edev; + const size_t xy_need = static_cast(M) * static_cast(H) * 2; + hipStream_t cst = static_cast(compute_q.handle); + + if (hipStreamWaitEvent(cst, tls.ev_e, 0) != hipSuccess) return false; + bool y_ok = + hipMemcpyPeerAsync(y_compute, compute_dev, tls.y, expert_dev, xy_need, cst) == hipSuccess; + if (!y_ok) { + std::vector tmp(xy_need); + Check(hipSetDevice(expert_dev), "y d2h"); + hipStream_t est = static_cast(tls.eq.handle); + if (hipStreamSynchronize(est) != hipSuccess) return false; + if (hipMemcpy(tmp.data(), tls.y, xy_need, hipMemcpyDeviceToHost) != hipSuccess) return false; + Check(hipSetDevice(compute_dev), "y h2d"); + if (hipMemcpyAsync(y_compute, tmp.data(), xy_need, hipMemcpyHostToDevice, cst) != hipSuccess) + return false; + } + tls.pending_M = 0; + if (hipSetDevice(compute_dev) != hipSuccess) return false; + return true; +} + +void PinGemma4Fp8ExpertHostCache(const Gemma4Fp8ExpertMats& ex) { + if (ex.host_pinned) return; + if (ex.cached_gu.empty() || ex.cached_dn.empty()) return; + const size_t gu_b = ex.cached_gu.size() * sizeof(uint16_t); + const size_t dn_b = ex.cached_dn.size() * sizeof(uint16_t); + hipError_t e1 = + hipHostRegister(const_cast(ex.cached_gu.data()), gu_b, hipHostRegisterDefault); + hipError_t e2 = diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-peerpipe-40-58.log b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-peerpipe-40-58.log new file mode 100644 index 000000000..319c0d480 --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-peerpipe-40-58.log @@ -0,0 +1,19 @@ +struct PeerSlot { + int cdev = -1, edev = -1; + int I = 0, H = 0, Mcap = 0; + int pending_M = 0; + vt::Queue eq{}; + bool eq_live = false; + void* x = nullptr; + void* y = nullptr; + void* gu_act = nullptr; + void* act = nullptr; + hipEvent_t ev_c = nullptr, ev_e = nullptr; +}; +struct PeerPipe { + PeerSlot s[2]; +}; +inline PeerPipe& PeerPipeTls() { + static thread_local PeerPipe pipe; + return pipe; +} diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-wrapper-950-1100.log b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-wrapper-950-1100.log new file mode 100644 index 000000000..f9804ca7a --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor-wrapper-950-1100.log @@ -0,0 +1,151 @@ + { + const hipError_t e = hipSetDevice(compute_dev); + if (e != hipSuccess) return fail("restore_cdev", e); + } + return true; +} + +// Prefill: run ExpertGeGLU for M rows on expert_dev (weights resident there). +// Peer only activations M×H bf16 — not multi-GB FP8 expert packs. +bool RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice(vt::Queue& compute_q, int expert_dev, + void* y_compute, const void* x_compute, + const void* fp8_gu, const void* s_gu, + const void* fp8_dn, const void* s_dn, int M, + int I, int H) { + const int compute_dev = compute_q.device.index; + if (expert_dev < 0 || compute_dev < 0 || !y_compute || !x_compute || !fp8_gu || !s_gu || + !fp8_dn || !s_dn || M <= 0 || I <= 0 || H <= 0) + return false; + + static const bool prefill_fp8_lt = [] { + if (const char* e = std::getenv("VT_GEMMA4_PREFILL_FP8_LT")) return e[0] == '1'; + return false; + }(); + static const bool gu_pair = [] { + const char* e = std::getenv("VT_GEMMA4_GU_INTERLEAVE"); + return e && e[0] == '1'; + }(); + + auto try_fp8_lt = [&](vt::Queue& q, void* y, const void* x, int m) -> bool { + if (!prefill_fp8_lt || m < 16) return false; + // gu_act[M, 2I] then GeluAndMul → act[M,I] then down → y[M,H] + // Scratch for gu_act/act via TLS-like statics on sticky device of q. + struct LtTls { + int dev = -1; + int I = 0, H = 0, Mcap = 0; + void* gu_act = nullptr; + void* act = nullptr; + }; + static thread_local LtTls lt; + const int dev = q.device.index; + const size_t gu_act_b = static_cast(m) * static_cast(2 * I) * 2; + const size_t act_b = static_cast(m) * static_cast(I) * 2; + if (lt.dev != dev || lt.I != I || lt.H != H || lt.Mcap < m || !lt.gu_act) { + auto freep = [&](void*& p) { + if (p) { + (void)hipFree(p); + p = nullptr; + } + }; + freep(lt.gu_act); + freep(lt.act); + Check(hipSetDevice(dev), "lt alloc"); + const int Mcap = std::max(m, 64); + if (hipMalloc(<.gu_act, static_cast(Mcap) * static_cast(2 * I) * 2) != + hipSuccess) + return false; + if (hipMalloc(<.act, static_cast(Mcap) * static_cast(I) * 2) != hipSuccess) + return false; + lt.dev = dev; + lt.I = I; + lt.H = H; + lt.Mcap = Mcap; + } + if (!vt::MatmulBTFp8ActWeightLt(q, lt.gu_act, x, fp8_gu, s_gu, m, 2 * I, H, 1.f, 0.f)) + return false; + vt::Tensor act = vt::Tensor::Contiguous(lt.act, vt::DType::kBF16, q.device, {m, I}); + vt::Tensor gu_act = vt::Tensor::Contiguous(lt.gu_act, vt::DType::kBF16, q.device, {m, 2 * I}); + if (gu_pair) + vt::rocm::GeluAndMulPairRocm(q, act, gu_act); + else + vt::GeluAndMul(q, act, gu_act); + if (!vt::MatmulBTFp8ActWeightLt(q, y, lt.act, fp8_dn, s_dn, m, H, I, 1.f, 0.f)) return false; + static thread_local int ok_log = 0; + if (ok_log++ < 3) { + std::fprintf(stderr, "INFO gemma4-moe prefill FP8×FP8 Lt path M=%d I=%d H=%d\n", m, I, H); + } + return true; + }; + + // Same device: caller should use local ExpertGeGLU; still support for safety. + if (expert_dev == compute_dev) { + if (try_fp8_lt(compute_q, y_compute, x_compute, M)) return true; + vt::Device dev = compute_q.device; + // Multi-slot sticky dequant (see PrefillDequantCache) + act scratch. + struct SameTls { + int dev = -1; + int I = 0, H = 0, Mcap = 0; + void* gu_act = nullptr; + void* act = nullptr; + }; + static thread_local SameTls tls_slots[2]; + SameTls& tls = tls_slots[(compute_dev == 1) ? 1 : 0]; + if (tls.dev != compute_dev || tls.I != I || tls.H != H || tls.Mcap < M || !tls.gu_act) { + auto freep = [&](void*& p) { + if (p) { + (void)hipFree(p); + p = nullptr; + } + }; + freep(tls.gu_act); + freep(tls.act); + Check(hipSetDevice(compute_dev), "same alloc"); + size_t gu_act_cap = static_cast(std::max(M, 64)) * static_cast(2 * I) * 2; + size_t act_cap = static_cast(std::max(M, 64)) * static_cast(I) * 2; + if (hipMalloc(&tls.gu_act, gu_act_cap) != hipSuccess) return false; + if (hipMalloc(&tls.act, act_cap) != hipSuccess) return false; + tls.dev = compute_dev; + tls.I = I; + tls.H = H; + tls.Mcap = std::max(M, 64); + } + auto& wslot = DequantCacheSlotFor(compute_dev); + void* gu_w = nullptr; + void* dn_w = nullptr; + int pin = -1; + { + std::lock_guard lk(wslot.cache.mu); + if (!wslot.cache.Ensure(compute_dev, I, H)) return false; + if (!wslot.cache.GetLocked(compute_q, fp8_gu, fp8_gu, s_gu, fp8_dn, s_dn, &gu_w, &dn_w, + &pin)) + return false; + } + vt::Tensor x = vt::Tensor::Contiguous(const_cast(x_compute), vt::DType::kBF16, dev, + {M, H}); + vt::Tensor gu_wt = vt::Tensor::Contiguous(gu_w, vt::DType::kBF16, dev, {2 * I, H}); + vt::Tensor gu_act = vt::Tensor::Contiguous(tls.gu_act, vt::DType::kBF16, dev, {M, 2 * I}); + vt::Tensor act = vt::Tensor::Contiguous(tls.act, vt::DType::kBF16, dev, {M, I}); + vt::MatmulBT(compute_q, gu_act, x, gu_wt); + if (gu_pair) + vt::rocm::GeluAndMulPairRocm(compute_q, act, gu_act); + else + vt::GeluAndMul(compute_q, act, gu_act); + vt::MatmulBTAlphaBeta(compute_q, y_compute, act.data, dn_w, M, H, I, 1.f, 0.f, + vt::DType::kBF16); + { + std::lock_guard lk(wslot.cache.mu); + wslot.cache.UnpinLocked(pin); + } + return true; + } + + // PEER path: dual-slot Launch+Finish (slot 0). Pipeline uses Launch/Finish directly. + if (!LaunchGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0, expert_dev, x_compute, fp8_gu, + s_gu, fp8_dn, s_dn, M, I, H)) + return false; + return FinishGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0, y_compute, M, H); +} + +bool LaunchGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, int expert_dev, + const void* x_compute, const void* fp8_gu, + const void* s_gu, const void* fp8_dn, const void* s_dn, diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper-donor.md b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor.md new file mode 100644 index 000000000..7297724b5 --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper-donor.md @@ -0,0 +1,38 @@ +# Donor evidence — #839 prefill peer helper + +Pinned **bytes**, not a dirty-tree HEAD. Implementation may use these slices as the +Launch/Finish / PeerPipe / cache **structure** donor. Product **must not** copy the +donor unpin-before-`ev_e` lifetime (research `64cb` stop-ship 3). See the spec. + +The slices sit beside the spec as `.log`, not under `.agents/evidence/`, because +`scripts/check-pr-size.py` classifies `.agents/specs/*.md` (`SPEC`) and +`.agents/specs/*.{patch,json,log}` (`SPEC_EVIDENCE`) and nothing else in +`.agents/`; an unclassified path makes that required check refuse the whole +change rather than the file. `.txt` is unclassified too, so the extension is +`.log`. Renaming does not touch the bytes and every SHA256 below is unchanged. + +| Field | Value | +|---|---| +| Donor tree | `/home/don/llms/vllm.cpp` | +| Donor git HEAD | `2bb4bd8a` (dirty; these slices are **uncommitted** on that tree) | +| File | `src/vt/rocm/rocm_gemma4_experts.hip` | +| `rocm-gemma4-prefill-peer-helper-donor-peerpipe-40-58.log` | lines 40–58 SHA256 `a40291b8db72fdbf41767516095ba358488afb92dbbb6cf90bf8c11697a3c005` | +| `rocm-gemma4-prefill-peer-helper-donor-dequant-cache-core-60-209.log` | lines 60–209 SHA256 `cdb4da131c5d4786bafac7b8f0601362c254409455773258961636ddfc43b937` (`PrefillDequantCacheSlots`, `Slot`, `FreeAll`, `Ensure`, `GetLocked`) | +| `rocm-gemma4-prefill-peer-helper-donor-dequant-cache-210-232.log` | lines 210–232 SHA256 `16de816ee053e9a0ae22d6eefb7ada9c811d99196f81a61047e5ac71796447c9` | +| `rocm-gemma4-prefill-peer-helper-donor-wrapper-950-1100.log` | lines 950–1100 SHA256 `e8c54fe8f415ca277bd5185f9e856dc98428030e241d6da819cd72cad107e098` | +| `rocm-gemma4-prefill-peer-helper-donor-launch-finish-1090-1310.log` | lines 1090–1310 SHA256 `ae1828d721e50a18b099bca4b08fc79cdc595410cf8d179f1f18744e0f6c1095` | +| Recipient | `origin/main` `3ce5a1dc` `rocm_gemma4_experts.hip:648` (monolithic) | +| Hanging compare | `vllm.cpp-bc64fa-r2` `1b1baf43` `:692` | +| Captured | 2026-08-14 | + +Known donor lifetime holes (do not productize as-is): + +- Launch unpins at ~1262–1265 immediately after enqueueing GEMMs and **before** `ev_e`; + Finish ~1274–1300 only enqueues `hipStreamWaitEvent` + copy and does **not** host-wait. +- Cache is process-wide (`DequantCacheSlotFor`, lines 220–231). +- `Ensure` (core 60–209) calls `FreeAll` on any device/I/H/nslots change; `FreeAll` + hipFrees and zeros **pinned** slots. Product must reject/defer reconfiguration + while any `pins > 0`. + +`sha256sum` of each slice file must match the table. Do not treat `2bb4bd8a` as a clean +donor commit. diff --git a/.agents/specs/rocm-gemma4-prefill-peer-helper.md b/.agents/specs/rocm-gemma4-prefill-peer-helper.md new file mode 100644 index 000000000..2b0f8806a --- /dev/null +++ b/.agents/specs/rocm-gemma4-prefill-peer-helper.md @@ -0,0 +1,220 @@ +# Spec: ROCm Gemma-4 prefill peer helper Launch/Finish + dequant-cache + +- **Issue:** https://github.com/mudler/vllm.cpp/issues/839 +- **Row slug:** `ROCM-GEMMA4-PREFILL-PEER-HELPER` — child of `BACKEND-ROCM` (#41). Separate from #697. +- **Worktree / branch (this unit only):** `/home/don/llms/vllm.cpp-prefill-peer` · `row/ROCM-GEMMA4-PREFILL-PEER-HELPER` +- **Base / recipient:** `origin/main` `3ce5a1dc` `rocm_gemma4_experts.hip:648` +- **Donor bytes:** `.agents/specs/rocm-gemma4-prefill-peer-helper-donor-*.log` slices hashed in [`rocm-gemma4-prefill-peer-helper-donor.md`](rocm-gemma4-prefill-peer-helper-donor.md) (dirty lab `/home/don/llms/vllm.cpp` HEAD `2bb4bd8a` **plus uncommitted**; HEAD is not a clean donor). +- **Implementer:** hermes-vllm. **Reviewer:** research. **Operator/smoke:** coord. +- **Git:** spec-only first (coord `25c9` / research `5071` / BLOCK `64cb`); impl after spec GREEN **on this same row branch**. Independent RED/GREEN from #837/#838. Ordered B then C is allowed; no all-at-once port. One PR per row. No shared `row/ROCM-GEMMA4-XDEV-MOE` landing history. +- **Depends on:** #837 is a separate hypothesis. Operator runs B then C on T=2029. Impl of C may assume B already landed **or** must still be A/B-able alone. +- **Supersedes for review:** `20332292` (BLOCK), preview `c4fbe6e9` (not spec-GREEN), and `231f38cf` (BLOCK `4954`). + +## Now + +`IMPLEMENTING` — ea9c: product RestoreComputeOrThrow no-op must RED via compile-and-run of the HIP body. + +## Adjacent upstream (not this row) + +- **#785**: `VT_ROCWMMA_OK` / #697 landing-order only. Do not touch `rocm_paged_attn.hip`. +- **#523 / #509**: custom keep-quant grouped GEMM + `rocm_moe_chain.hip`. No GetBlas/Launch/Finish overlap; docs-only if FEATURES/USAGE touch. +- **#834**: router-lookahead prefetch — adjacent cache policy, not this pin lifetime. + +**Not a confirmed fix.** Hypothesis (C) only. `9772` does not isolate Launch/Finish vs dequant-cache vs GetBlas. Repeated invocation is **observed** to wedge after matched BEGIN/END. Cause (single-slot TLS vs cache eviction vs GetBlas destroy) is **unconfirmed**. Do not say repeated invocation "exhausts" single-slot state. + +This package stays one row because the lab wrapper is one ownership unit; tests must still fail independently if pin-lifetime or Launch/Finish pairing is broken. + +## Upstream / source of the port + +No vLLM Python equivalent. Source is the pinned donor slices for **structure**. Product lifetime is **stricter** than the donor (see Design). + +| Tree | Symbol | Shape | +|---|---|---| +| `origin/main` `3ce5a1dc` | `RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice` `:648` | one function; same-dev `static thread_local SameTls tls` + sticky; peer `static thread_local Tls tls` + inline copy/GEMM/`ev_c`/`ev_e` | +| hanging `vllm.cpp-bc64fa-r2` `1b1baf43` | `:692` | same monolithic body | +| donor slices | wrapper `:959-1095`; `Launch…` `:1098`; `Finish…` `:1274`; `PeerPipeTls` `:55`; cache core `:60-209`; `DequantCacheSlotFor` `:225` | wrapper Launch(slot0)→Finish; `PrefillDequantCache` Slot/FreeAll/Ensure/GetLocked; same-dev `SameTls tls_slots[2]` + cache pin/unpin | + +`kPeerPipe` in lab is **default OFF**. Even OFF, the wrapper + dequant-cache lifetime is the delta (`713f`). + +## Symptom this row owns + +T≥64 / T=2029 prefill-batch peer path. Coord `9772`: every individual `RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice` returns; 274 BEGIN/END match; then idle + `kfd_wait_on_events`. Accumulation **class**, not a proven single-op hang. p42k-critical. Cause unconfirmed. + +## Scope + +Port the **resource-managed** lab shape onto main's helper, then close the donor pin hole: + +1. Keep the public symbol `RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice` as a wrapper. +2. Same-dev: `SameTls tls_slots[2]`; weights via `DequantCacheSlotFor(compute_dev)` pin → GEMM → **retire → unpin**. No sticky `const void*` in the act TLS. +3. Peer: `LaunchGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, slot=0, …)` then `FinishGemma4Fp8ExpertGeGLUPrefillPeer(...)`. +4. `PeerPipeTls().s[slot]` holds x/y/events/expert queue **and the live cache pin** for that slot. +5. `DequantCacheSlotFor(expert_dev)` pin around expert-side dequant+GEMM; unpin only after GPU consumption retires. +6. `kPeerPipe` / any overlap env **default OFF**. Pipe-on is out of this slice. + +Caller in `gemma4_moe.cpp` prefill-batch peer-act chunk loop stays the public wrapper. Do not change gather/scatter or `kPrefillBatchMinT`. + +## Out of scope + +- Enabling peer-pipe overlap as product default. +- `VT_GEMMA4_PREFILL_FP8_LT` (KEEP-rejected) and `VT_GEMMA4_GU_INTERLEAVE`. +- GetBlas body (#837) — but impl of this row must not ship without #837 landed or bundled as a **separate** immutable head. +- Indexed T<63 (#838). +- `#697` / `rocm_paged_attn.hip`. +- Diagnostic `STAGE_SYNC` / `PREFILL_TRACE` / M1 drains / `hipSetDevice` added only for prints. +- Transplant of dirty-lab layer-split / FIFO / resident-pack extras. +- Copying donor unpin-before-`ev_e`. + +## Design + +`kPeerPipe` OFF means Launch and Finish are back-to-back on slot 0 — **same event order as today's wait-then-y**, plus slot-scoped scratch. + +### Donor hole (must not ship) + +Donor Launch (`rocm-gemma4-prefill-peer-helper-donor-launch-finish-1090-1310.log`) unpins at ~1262–1265 immediately after enqueueing GEMMs and **before** `hipEventRecord(tls.ev_e)`. Finish ~1274–1300 enqueues `hipStreamWaitEvent(cst, tls.ev_e)` + peer copy and does **not** host-wait `ev_e`. `DequantCacheSlotFor` is process-wide (220–231), so another worker/stream may evict/rewrite a zero-pin slot while the first expert stream still reads it. "Unpin after Finish" is still insufficient unless Finish proves `ev_e` complete. + +Donor `Ensure` (`rocm-gemma4-prefill-peer-helper-donor-dequant-cache-core-60-209.log`) calls `FreeAll` whenever device/I/H/`nslots` disagree with the request. `FreeAll` `hipFree`s every slot and zeros `pins`, including live pins. Product must not copy that. + +### Product ownership (required) + +Persist the pin in `PeerSlot`: + +``` +PeerSlot { + ... existing x/y/eq/ev_c/ev_e/pending_M/edev ... + int cache_pin = -1; // slot index inside DequantCache + int cache_dev = -1; // device whose DequantCacheSlotFor owns it +} +``` + +Rules: + +1. Launch pins under the same mutex scope as donor `GetLocked` (lock → GetLocked → unlock → GEMM). Store `{cache_pin, cache_dev}` on the `PeerSlot` **before** any GEMM enqueue. Do **not** unpin in Launch. +2. Finish may enqueue `hipStreamWaitEvent` + y-copy, then **must host-wait** `ev_e` (or `hipStreamSynchronize` on the expert stream that recorded it) **before** unpin. `hipStreamWaitEvent` on the compute stream is not a host-side retirement proof. +3. Unpin under lock via `UnpinLocked` only after that host-wait succeeds. Then clear `cache_pin = -1`. +4. Same-dev enqueue-only path: same rule. Record a done-event on the GEMM stream; host-wait it before unpin. "Enqueue then unpin" is forbidden even when compute_dev == expert_dev, unless a written concurrency argument proves no other worker can `GetLocked`/`evict` that slot — default is: **no such argument**, so host-wait. +5. Every Launch/Finish error path after a successful `GetLocked` must retire GPU work **before** unpin: + - If `ev_e` was successfully recorded: host-wait `ev_e` (or `hipStreamSynchronize` on the stream that recorded it), then unpin. + - If pin is live and `ev_e` was **not** recorded (fail after pin, after dequant/GEMM enqueue, or `hipEventRecord(ev_e)` itself fails): **`hipStreamSynchronize` the expert stream** (same-dev: the compute/GEMM stream) **then** unpin. "Wait any recorded `ev_e` if valid" is not a retirement proof when the event does not exist. + - Then restore compute device. No live expert event left. Pin count → 0 except failed-retirement quarantine (pin stays until later observed retire or fatal teardown). +6. Mutex scope still mirrors donor for Get/Unpin (`GetLocked` under lock, GEMM outside, `UnpinLocked` under lock). Do not hold the mutex across `MatmulBT`. +7. **`Ensure` must not reconfigure while any pin is live.** If `dev/I/H/nslots` already match, return true. If they differ and any `slots[i].pins > 0`, return false (reject; do not `FreeAll`). Only call `FreeAll` when pin-count is 0 across the cache. Do not copy donor `Ensure`→`FreeAll` on a live pin. Slot-count comes from `PrefillDequantCacheSlots()` (donor default 1). + +### Lifetime invariants (mutation-proven) + +- every successful Launch has exactly one Finish; +- Launch fail unpins (or never pinned) after current-stream retire; leftover prior ev_e is never the rollback target; failed retire quarantines; +- Finish restores compute device before return; +- `ev_c` is recorded on compute stream, `ev_e` on expert stream; no cross-device event record; +- cache pin count returns to 0 on success and on every error return where retirement was observed; failed retirement is the sole zero-pin exception (quarantine); +- pin remains >0 from GetLocked until host-observed retirement (`ev_e` complete, or expert/compute stream sync on the pre-record error path); +- a second worker cannot obtain a rewrite/evict of a still-pinned slot; +- `Ensure` never `FreeAll`s a cache that has any `pins > 0`. + +## Risks + +- Host-wait on every Finish adds latency vs donor enqueue-only unpin. Accept for correctness. Pipe-on overlap is a later spec. +- Without #837, Finish-side `MatmulBT` can still destroy the other GPU's hipBLAS handle. +- Slot 0 only while pipe is off. Do not "simplify" by keeping monolithic TLS. + +## Tests + +Host-only: + +1. Source invariant: product wrapper contains `LaunchGemma4Fp8ExpertGeGLUPrefillPeer` and `FinishGemma4Fp8ExpertGeGLUPrefillPeer`; does not keep a peer `static thread_local Tls tls` as the only storage. +2. Source invariant: `PeerSlot` stores `cache_pin`; Launch does not call `UnpinLocked` before `ev_e` record; Finish host-waits before unpin; `Ensure` returns false when any pin is live instead of `FreeAll`. +3. Pairing: every Launch path that returns true has a Finish; fail-Launch must unpin. +4. `kPeerPipe` default is off (env unset → no overlapping slot 1). + +RED: restore monolithic peer TLS → invariant 1 fails. RED: restore donor unpin-before-`ev_e` → invariant 2 fails. + +Host or fake-cache mutations (required): + +- concurrent eviction: worker B `GetLocked`/rewrite of a zero-pin slot while worker A's expert stream still "reads" (fake outstanding GEMM). Must RED if unpin happened before retirement. +- fail after pin, **after work enqueue** (dequant or GEMM queued), before event record → expert-stream (same-dev: compute-stream) sync then unpin; pin count 0; no live event. A mutation that only fails immediately after pin (no enqueue) does **not** satisfy this case. +- fail at `hipEventRecord(ev_e)` after work enqueue → same stream-sync fallback before unpin. +- fail after event record, before wait → pin count 0 after cleanup, event retired or destroyed safely. +- fail after wait, before copy → pin count 0. +- fail after copy, before unpin → cleanup still unpins (idempotent). +- concurrent `Ensure` with a new I/H or slot-count while any pin is live → `Ensure` returns false, no `FreeAll`, pinned buffers remain. RED if donor `Ensure`→`FreeAll` is restored. + +GPU (coord, after #837): + +- T=2029, `VT_GEMMA4_PREFILL_TRACE=1`, `VT_GEMMA4_ATTN_STAGE_SYNC` **unset**, `MAX_BATCHED_TOKENS=8192`, `MAX_MODEL_LEN=65536`, `NUM_BLOCKS=2048`, `PEER_ACT=1`: matched helper BEGIN/END **and** HTTP body. +- T=1 decode + Paris/arith unchanged. + +## Gates + +- Host invariants GREEN without GPU. +- No FP8 Lt / GU_INTERLEAVE / pipe-default-on in the impl. +- Operator A/B (`5071`): this is **C**. After B, smoke **T=2029** independently (console+wchan on failure). p42k only after the smallest passing set of A/B/C. +- Default behavior outside Gemma-4 FP8 xdev prefill-batch peer-act is unchanged. +- `#697` files untouched. + +## Stop conditions + +- Research requires pipe-on default or FP8 Lt — new spec. +- Lab GPU / p42k without coord. +- Impl that edits `rocm_paged_attn.hip`. +- Impl that copies donor unpin-before-`ev_e`. +- Landing this row on a shared branch with #837/#838. + +## Evidence + +Bus: `0384`, `713f`, `9772`, `25c9`, `5071`, `64cb`, `4954`. Donor bytes hashed in [`rocm-gemma4-prefill-peer-helper-donor.md`](rocm-gemma4-prefill-peer-helper-donor.md) (cache core 60–209 required). + +Now: `IMPLEMENTING` — ea9c product RestoreComputeOrThrow no-op compile-and-run RED. `a007ec40` is not a review target. + +## Owed + +Owned by this row (`ROCM-GEMMA4-PREFILL-PEER-HELPER`) and tracked by +[#839](https://github.com/mudler/vllm.cpp/issues/839), which stays open until each +is discharged. Named here rather than left to be discovered, per AGENTS.md +"Nothing lands dead". + +- **The host simulator is a second implementation of Launch/Finish.** + `include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h` carries `HostAlloc`, + `PrefillDequantCacheHost`, `HostRetireThenUnpin`, `HostLaunch` and `HostFinish` + (~250 of its lines), compiled into every HIP build. `HostLaunch`/`HostFinish` + are hand-written analogues of the product `LaunchGemma4Fp8ExpertGeGLUPrefillPeer` + and `FinishGemma4Fp8ExpertGeGLUPrefillPeer` in + `src/vt/rocm/rocm_gemma4_experts.hip`, so the lifetime cases in + `tests/vt/test_ops_gemma4_prefill_peer.cpp` exercise the analogue and not the + product. What binds the two today is the source-slice mutation gates in the + same file (`ProductFinishRetiresCstBeforeReuse`, `CompileFinishCatchTu`, + `CompileAndRunProductRestore`), which assert against the product TEXT. Owed: + move the `Host*` half under `tests/`, keeping `PrefillPeerLife`, + `ChoosePrefillRetire`, `PrefillDequantCacheT`, `ComputeDevGuard`, + `RestoreFailed`, `SameDevLife`, `OutputCopyGate`, `SlotReusable` and + `PublishThenRestoreOrThrow` in the header because the `.hip` reaches them. + Deliberately not done in this change: only a ROCm box compiles the one + translation unit that consumes the header, so the move cannot be verified + where this repair was made. + +- **The blocking retirement has no measurement — protocol frozen to Researcher + ca41. Do NOT run the historical async-unpin product arm.** That arm is the + lifetime bug; using it as a “before” product binary can unpin storage while + queued GEMMs still consume it. Safe attribution only: + 1. Env-gated measurement-only counters around each product + `hipEventSynchronize(ev_e)` and `hipStreamSynchronize(cst)`: call count, + host-block microseconds by kind, and pre-sync + `hipEventQuery`/`hipStreamQuery` ready-vs-not-ready. No change to + ordering/lifetime. + 2. Frozen prefill smoke T=2029 on 2×R9700: uninstrumented current-head + throughput plus instrumented attribution as separate legs. Same + binary/head/recipe; ≥3 warmups + ≥5 steady samples. Report prefill + tok/s/TTFT, total request time, wait calls, cumulative wait-host-us, + us/call, ready fraction, wait-time/request-time ratio. Instrumented + throughput is diagnostic, not canonical. + 3. If a literal before/after remains mandatory: a **separate synthetic HIP + microbench** whose buffers stay live until one final stream sync. A waits + per expert like product; B enqueues the same sequence then waits once. + Same op count/streams/bytes. Label as overlap upper bound, not product + throughput. Never disable retirement/unpin in the model path. Never infer + `before = elapsed - wait_us` as measured throughput. + 4. Isolated `:8012`; never `:8010`; idle window; teardown. Record raw logs + + HEAD/binary/recipe in `docs/BENCHMARKS.md` + `.agents/benchmark-record.md`. + Expensive → next-row; the correctness wait stays. + +- **`PeerSlot s[1]` SOURCE FIXED (Researcher ca41).** No slot-1 overlap wiring. + Launch/Finish fail-closed on `slot != 0`. Wrapper still hardcodes `/*slot=*/0`. + GPU not required for this item. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 4359c3972..90a631701 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -29,6 +29,7 @@ | **DwarfStar** | DeepSeek-V4-Flash GGUF, GB10 | **beats ds4, 1.144x** (18.69 vs 16.33 tok/s, byte-exact, default config) | n/a, GGUF peer | | **vLLM** | Kimi-Linear-48B-A3B, GB10 | no binding number: the published checkpoint is tiktoken-only, so it cannot drive the warm-server harness | golden 122/128, near-tie profile | | **Muse Glimmer 30B (#268)** | no vLLM denominator (pin lacks `muse_glimmer`); SECONDARY llama.cpp, same GGUF, idle GB10, after the [#391 fix](../.agents/specs/cpu-decode-barrier-and-attn-dispatch.md) | **vLLM axis is an OPEN GAP.** vs llama.cpp: in128 prefill **1.023x** (was 0.878x); in512 decode **0.194x** (3.41x), prefill 0.175x flat. Denominator stock `7044859`, SUPERSEDED (#1003) | coherent, NOT token-exact | +| **#1047 prefill-peer Finish barriers (item 3)** | Gemma-4-26B-A4B-it-fp8, dual R9700, isolated :8012, T=2029, PEER_ACT=1 | **Attribution only, not a ship claim.** Two Finish success-path barriers cost **2.55%** (1122.10 vs 1094.24 tok/s; 46.05 ms/req) vs retained RetirePinThenUnpin wait. [record](../.agents/benchmark-record.md) | ident-equal 4/4 | Reading the ratios: throughput is ours/reference, latency is reference/ours, so **1.0 or higher is a win** everywhere on this page. Which architecture each number @@ -511,6 +512,7 @@ built on it rather than keeping the flattering one. | Vulkan vs llama.cpp Vulkan ([`BENCH-VK-LLAMA`](../benchmarks/demo/vulkan_27b_llamacpp.json)) | 25 NATIVE (+8 GDN). **27B prefill 21.5x, a SELF-ratio not a llama.cpp one**; decode **4.36 vs 4.35, MET**, denominator SUPERSEDED (7 clean legs). Smart barriers skip 19.8%/tok, GPU -1.09 ms; e2e 8/12, unresolved. OFF. | `VK-C` coopmat A/B on Thor (`VT_VULKAN_COOPMAT=0` A/Bs it): **11.1x-32.9x** vs our UNTILED scalar kernel, not a competent GEMM. `VK-E`: llama.cpp `-DGGML_VULKAN=ON` at `237ad9b96` on dgx, SUPERSEDED, same GGUF, 3 columns | | ROCm (`BACKEND-GATE-ROCM-VLLM` / `-SGLANG`) | **PENDING: no binding throughput number.** The directional `d=128` decode row below is not one. Runtime-green on 5 gfx archs. Gemma-3 is 48/48 exact vs two vLLM-ROCm oracles; Qwen3.5-0.8B correctness remains open | Same model, quantization, request shape and cache policy vs pinned vLLM-ROCm on one idle AMD host. Add equivalent SGLang; close correctness first ([#41](https://github.com/mudler/vllm.cpp/issues/41)) | | ROCm `d=128` decode arm (`BACKEND-ROCM`, [#382](https://github.com/mudler/vllm.cpp/issues/382)) | **DIRECTIONAL, not binding.** gfx1200, both sides in the pinned oracle container, 1024/128 c1, 8 prompts, 3 reps: TPOT 42.40 -> **11.67 ms** with `VT_ATTN_DECODE_D128=1`; vLLM `555967922` 6.68 ms, so **6.35x -> 1.75x** | Harnesses differ (oracle over HTTP, ours in-process) and no same-tool per-call trace exists, so [#488](https://github.com/mudler/vllm.cpp/issues/488) stays open. Owed: decode-windowed `rocprofv3` both sides | +| Gemma4 prefill-peer Finish-barrier cost (#1047 item 3) | **Attribution GREEN at T=2029, not a product ship number.** A (wait-only BEFORE) 1122.10 tok/s vs B (AFTER) 1094.24; +2.55% / 46.05 ms/req. Wait deletion is not authorized to land. | T=19 not run. Overlap/async retirement unmeasured. Detail: [benchmark-record](../.agents/benchmark-record.md) | | Tenstorrent Blackhole (`BACKEND-TENSTORRENT`) | **NOT APPLICABLE (speed).** Correctness: OPT-125m STRICT 6/6 e2e on real hardware. Qwen3-0.6B has a device-specific golden and short 4-token warm smoke (~0.28 tok/s), not a completed speed run | Full 16x16 Qwen3 gate, then device-resident tensors + `ttnn::sdpa_decode` before any performance comparison. [Spec](../.agents/specs/tenstorrent-backend.md) | | Mistral-7B-v0.3 on Tenstorrent (`BACKEND-TENSTORRENT-MISTRAL`) | **PENDING (speed).** First data point: 4.26 tok/s warm, batch 1, 32 tok, single run on a P150. Not a gate, not reproduced. No vLLM ratio exists or can (no TT backend). Correctness 16/16 | Reproduce idle with a same-binary A/B before quoting. [Record](../.agents/benchmark-record.md), [spec](../.agents/specs/tenstorrent-mistral.md) | | Host-free decode graph (`BACKEND-TENSTORRENT-HOST-FREE-FORWARD`) | **PENDING (operator gate).** Implementer P150 run (2026-08-16): Qwen3-0.6B, 80 tok, 5.8x vs eager, 79 replays, 22/22 vs copy baseline. Env-gated. No vLLM ratio | Operator rerun, then full-engine `test_qwen3_paged_engine`. [Spec](../.agents/specs/tenstorrent-host-free-forward.md) | diff --git a/docs/STATUS.md b/docs/STATUS.md index 82ac0e8f9..c639ce7ab 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -572,7 +572,7 @@ output-dtype dispatch fix; gfx1201 Gemma-4 FP8 MoE is contributor-measured on 2x R9700 and CPU-link-verified our side; a `head_dim=128` decode arm lands opt-in behind `VT_ATTN_DECODE_D128`, default OFF, which moves gfx1200 per-token decode from 6.35x to 1.75x slower than the pinned vLLM oracle on one shape, a -directional figure that leaves the ROCm throughput axis PENDING; +directional figure that leaves the ROCm throughput axis PENDING; #1047 item 3 on gfx1201 Gemma-4 FP8 prefill-peer is attribution only (T=2029: deleting two Finish success-path barriers is +2.55% / 46.05 ms/req vs the retained RetirePinThenUnpin wait, not a license to land the deletion and not a ROCm throughput pin; [BENCHMARKS](BENCHMARKS.md)); [guide](ROCM.md)), inference-time CPU weight offload (`ENG-WEIGHT-OFFLOAD` ACTIVE; the config surface landed W0a (the backend enum, both sub-configs, the validator's two errors and three warnings, and the dot-anchored segment match), diff --git a/docs/USAGE.md b/docs/USAGE.md index cff3a4c99..527854b15 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -4429,6 +4429,7 @@ GetBlas keeps two per-thread hipBLAS handles (`tls_slots[2]`, device 1 → slot 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. +Prefill peer (#839) unpins dequant cache only after observed retirement; a failed fill/ready lease is retired with RetireFillLocked after the producer stream sync (never under cache.mu); restore-fail after publish retires before rethrow; failed retire quarantines the pin. This path 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). @@ -5129,3 +5130,7 @@ length is the one the request's duration implies, and that it is **real audio** non-zero, unclipped, non-constant, and with two channels that differ (the stereo fold is a contiguous split of the 128 latent channels, and an interleave produces a correctly shaped, correctly ranged, wrong song). + +Gemma-4 FP8 xdev prefill (`RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice`) is a +Launch/Finish wrapper: cache pins stay live until host-observed `ev_e` retirement. +Peer-pipe overlap stays off (slot 0 only). diff --git a/include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h b/include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h new file mode 100644 index 000000000..729d6ad4a --- /dev/null +++ b/include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h @@ -0,0 +1,624 @@ +// #839 shared prefill-peer lifetime. Product HIP and host tests execute these +// transitions. Failed retirement quarantines the pin (sole zero-pin exception). +#pragma once + +#include +#include +#include +#include +#include +#include + +namespace vt::rocm { + +inline int PrefillDequantCacheSlots() { return 1; } +inline constexpr int kPrefillDequantCacheMaxSlots = 128; + +enum class PrefillRetireOutcome { None, Unpinned, Quarantined }; +enum class PrefillRetireTarget { None, ComputeStream, ExpertStream, RecordedEvent }; + +// Finish output copy on a specific compute stream. Reuse/reconfigure of tls.y +// is forbidden until that stream is host-observed retired. +struct OutputCopyGate { + int copy_stream = -1; + bool pending = false; + void Enqueue(int sid) { + copy_stream = sid; + pending = true; + } + bool Retire(int sid) { + if (!pending) return true; + if (sid != copy_stream) return false; + pending = false; + copy_stream = -1; + return true; + } + bool CanReuseScratch() const { return !pending; } +}; + +struct PrefillPeerLife { + int cache_pin = -1; + int cache_dev = -1; + int pending_M = 0; + bool ev_e_recorded = false; + bool this_gen_ev_e = false; + bool work_enqueued = false; + bool work_on_compute = false; + bool work_on_expert = false; + bool rollback_armed = false; + bool quarantined = false; + bool compute_restored = false; + bool fill_lease = false; + OutputCopyGate output_copy{}; + + // Arm after the first successful current-generation enqueue. + void ArmRollback() { + rollback_armed = true; + work_enqueued = true; + this_gen_ev_e = false; + } + void MarkComputeWork() { work_on_compute = true; } + void MarkExpertWork() { work_on_expert = true; } + void MarkThisGenEvent() { + this_gen_ev_e = true; + ev_e_recorded = true; + rollback_armed = false; + } + void OnSuccessfulRetire() { + cache_pin = -1; + cache_dev = -1; + ev_e_recorded = false; + this_gen_ev_e = false; + rollback_armed = false; + work_on_compute = false; + work_on_expert = false; + fill_lease = false; + } +}; + +// Map any slot-like object that carries the product flags. +template +PrefillPeerLife LifeFromSlot(const Slot& tls) { + PrefillPeerLife life; + life.cache_pin = tls.cache_pin; + life.cache_dev = tls.cache_dev; + life.this_gen_ev_e = tls.this_gen_ev_e; + life.ev_e_recorded = tls.ev_e_recorded; + life.rollback_armed = tls.rollback_armed; + life.quarantined = tls.quarantined; + life.work_on_compute = tls.work_on_compute; + life.work_on_expert = tls.work_on_expert; + life.fill_lease = tls.fill_lease; + return life; +} + +inline bool SlotReusable(bool pending, bool rollback_armed, bool quarantined) { + return !pending && !rollback_armed && !quarantined; +} + +// this_gen ev_e wins. Else expert-stream work, else compute-stream work. +// Never treat a leftover ev_e_recorded as the current rollback target. +inline PrefillRetireTarget ChoosePrefillRetire(const PrefillPeerLife& life) { + if (life.this_gen_ev_e) return PrefillRetireTarget::RecordedEvent; + if (life.work_on_expert) return PrefillRetireTarget::ExpertStream; + if (life.work_on_compute || life.rollback_armed || life.cache_pin >= 0) + return PrefillRetireTarget::ComputeStream; + return PrefillRetireTarget::None; +} + +struct RestoreFailed : std::runtime_error { + RestoreFailed() : std::runtime_error("prefill-peer compute device restore failed") {} +}; + +template +struct ComputeDevGuard { + int dev = -1; + SetDev set{}; + bool done = false; + + ComputeDevGuard(int d, SetDev s) : dev(d), set(std::move(s)) {} + ComputeDevGuard(const ComputeDevGuard&) = delete; + ComputeDevGuard& operator=(const ComputeDevGuard&) = delete; + + void RestoreOrThrow() { + if (done) return; + const bool ok = set(dev); + done = true; // failure is thrown; dtor must not terminate on the same attempt + if (!ok) throw RestoreFailed{}; + } + + // Failed restore is fatal even during exception unwind (c2ae). + ~ComputeDevGuard() noexcept { + if (done) return; + if (!set(dev)) std::terminate(); + done = true; + } +}; + +struct SameDevLife { + int cache_pin = -1; + int cache_dev = -1; + bool quarantined = false; + bool CanEnter() const { return !quarantined; } + bool CanReconfigure() const { return !quarantined && cache_pin < 0; } + void PersistPin(int pin, int dev) { + cache_pin = pin; + cache_dev = dev; + } + void Quarantine(int pin, int dev) { + cache_pin = pin; + cache_dev = dev; + quarantined = true; + } + void ClearPin() { + cache_pin = -1; + cache_dev = -1; + quarantined = false; + } +}; + +template +struct PrefillDequantCacheT { + struct Slot { + const void* key = nullptr; + void* gu = nullptr; + void* dn = nullptr; + uint64_t lru = 0; + int pins = 0; + bool ready = false; + bool filling = false; + bool fill_failed = false; + typename Hooks::Event ready_ev{}; + }; + + int dev = -1; + int I = 0, H = 0; + int nslots = 0; + uint64_t clock = 0; + uint64_t hits = 0, misses = 0; + Slot slots[kPrefillDequantCacheMaxSlots]{}; + std::mutex mu; + Hooks hooks{}; + + void FreeAll() { + for (int i = 0; i < kPrefillDequantCacheMaxSlots; ++i) { + if (slots[i].gu) { + hooks.Free(slots[i].gu); + slots[i].gu = nullptr; + } + if (slots[i].dn) { + hooks.Free(slots[i].dn); + slots[i].dn = nullptr; + } + hooks.DestroyEvent(slots[i].ready_ev); + slots[i] = Slot{}; + } + nslots = 0; + dev = -1; + I = H = 0; + } + + int LivePins() const { + int n = 0; + for (int i = 0; i < kPrefillDequantCacheMaxSlots; ++i) n += slots[i].pins; + return n; + } + + int LiveAllocs() const { + int n = 0; + for (int i = 0; i < kPrefillDequantCacheMaxSlots; ++i) { + if (slots[i].gu) ++n; + if (slots[i].dn) ++n; + } + return n; + } + + bool SlotNonReusable(int i) const { + return slots[i].filling || slots[i].fill_failed || slots[i].pins > 0; + } + + bool Ensure(int device, int i_dim, int h_dim) { + const int want = PrefillDequantCacheSlots(); + if (dev == device && I == i_dim && H == h_dim && nslots == want) return true; + if (LivePins() > 0) return false; + FreeAll(); + if (device < 0 || i_dim <= 0 || h_dim <= 0 || want <= 0) return false; + if (!hooks.SetDevice(device)) return false; + const size_t gu_b = static_cast(2 * i_dim) * static_cast(h_dim) * 2; + const size_t dn_b = static_cast(h_dim) * static_cast(i_dim) * 2; + for (int i = 0; i < want; ++i) { + slots[i].gu = hooks.Malloc(gu_b); + if (!slots[i].gu) { + FreeAll(); + return false; + } + slots[i].dn = hooks.Malloc(dn_b); + if (!slots[i].dn) { + FreeAll(); + return false; + } + } + dev = device; + I = i_dim; + H = h_dim; + nslots = want; + clock = 0; + return true; + } + + // pin_out is written as soon as a fill lease or hit pin is taken, including + // failed Fill/RecordReady — caller must persist/retire that pin. + bool GetLocked(const void* key, void** gu_out, void** dn_out, int* pin_out) { + if (!key || !gu_out || !dn_out || nslots <= 0) return false; + int hit = -1; + int victim = -1; + uint64_t oldest = UINT64_MAX; + for (int i = 0; i < nslots; ++i) { + if (slots[i].key == key && slots[i].ready && !slots[i].filling && !slots[i].fill_failed) { + hit = i; + break; + } + if (SlotNonReusable(i)) continue; + if (slots[i].key == nullptr) { + if (victim < 0) victim = i; + } else if (slots[i].lru < oldest) { + oldest = slots[i].lru; + if (victim < 0 || slots[victim].key != nullptr) victim = i; + } + } + if (hit >= 0) { + if (!hooks.WaitReady(slots[hit].ready_ev)) return false; + ++hits; + slots[hit].lru = ++clock; + slots[hit].pins++; + *gu_out = slots[hit].gu; + *dn_out = slots[hit].dn; + if (pin_out) *pin_out = hit; + return true; + } + if (victim < 0) { + for (int i = 0; i < nslots; ++i) { + if (!SlotNonReusable(i)) { + victim = i; + break; + } + } + if (victim < 0) return false; + } + Slot& s = slots[victim]; + s.key = nullptr; // unpublish before fill + s.ready = false; + s.filling = true; + s.fill_failed = false; + s.pins++; // fill lease from first enqueue + if (pin_out) *pin_out = victim; + if (!hooks.Fill(s.gu, s.dn)) { + s.fill_failed = true; + return false; + } + if (!hooks.RecordReady(s.ready_ev)) { + s.fill_failed = true; + return false; + } + s.ready = true; + s.key = key; + s.filling = false; + ++misses; + s.lru = ++clock; + *gu_out = s.gu; + *dn_out = s.dn; + return true; + } + + void UnpinLocked(int idx) { + if (idx < 0 || idx >= nslots) return; + if (slots[idx].pins > 0) slots[idx].pins--; + } + + bool IsFillLease(int idx) const { + if (idx < 0 || idx >= nslots) return false; + return slots[idx].filling || slots[idx].fill_failed; + } + + // After a successful owning-stream retire: fill-lease uses RetireFillLocked, + // ready pin uses UnpinLocked. Caller must not hold mu around the stream sync. + bool ReleaseObservedPinLocked(int idx, bool fill_lease) { + if (idx < 0 || idx >= nslots) return true; + if (fill_lease || IsFillLease(idx)) return RetireFillLocked(idx, true); + UnpinLocked(idx); + return true; + } + + // Producer-stream retirement of a failed fill lease. + bool RetireFillLocked(int idx, bool retire_ok) { + if (idx < 0 || idx >= nslots) return false; + if (!retire_ok) { + slots[idx].fill_failed = true; + return false; + } + slots[idx].filling = false; + slots[idx].fill_failed = false; + if (slots[idx].pins > 0) slots[idx].pins--; + return true; + } +}; + +struct HostAlloc { + int next = 1; + int mallocs = 0; + int frees = 0; + int fill_calls = 0; + int fail_malloc_at = -1; + bool fail_record = false; + bool fail_wait = false; + bool fail_set_device = false; + bool fail_fill = false; + struct Event { + bool recorded = false; + }; + bool SetDevice(int) { return !fail_set_device; } + void* Malloc(size_t) { + ++mallocs; + if (fail_malloc_at >= 0 && mallocs == fail_malloc_at) return nullptr; + return reinterpret_cast(static_cast(++next)); + } + void Free(void* p) { + if (p) ++frees; + } + void DestroyEvent(Event& e) { e.recorded = false; } + bool Fill(void*, void*) { + ++fill_calls; + return !fail_fill; + } + bool RecordReady(Event& e) { + if (fail_record) return false; + e.recorded = true; + return true; + } + bool WaitReady(Event& e) { return e.recorded && !fail_wait; } +}; + +using PrefillDequantCacheHost = PrefillDequantCacheT; + +inline PrefillRetireOutcome RetirePinIfObserved(PrefillDequantCacheHost& cache, int& cache_pin, + bool retire_ok, bool fill_lease = false) { + if (cache_pin < 0) return PrefillRetireOutcome::None; + if (!retire_ok) return PrefillRetireOutcome::Quarantined; + std::lock_guard lk(cache.mu); + if (!cache.ReleaseObservedPinLocked(cache_pin, fill_lease)) return PrefillRetireOutcome::Quarantined; + cache_pin = -1; + return PrefillRetireOutcome::Unpinned; +} + +// Product publish-then-restore: retire/quarantine before rethrowing RestoreFailed. +template +void PublishThenRestoreOrThrow(Slot& slot, int M, Restore&& restore, Retire&& retire) { + slot.pending_M = M; + slot.rollback_armed = false; + try { + restore(); + } catch (const RestoreFailed&) { + (void)retire(); + slot.pending_M = 0; + throw; + } +} + +enum class PrefillPeerFailAt { + None, + AfterFirstEnqueue, + AfterPinEnqueue, + RecordEvE, + AfterRecord, + AfterWait, + AfterCopy, + AfterSameDevAcquire, + AfterSameDevGemm, +}; + +using PrefillPeerSlotHost = PrefillPeerLife; + +inline bool HostRetireThenUnpin(PrefillDequantCacheHost& cache, PrefillPeerLife& slot, + bool retire_ok) { + const auto target = ChoosePrefillRetire(slot); + if (target == PrefillRetireTarget::None && slot.cache_pin < 0) return true; + if (!retire_ok) { + slot.quarantined = true; + return false; + } + const auto out = RetirePinIfObserved(cache, slot.cache_pin, /*retire_ok=*/true, slot.fill_lease); + if (out == PrefillRetireOutcome::Unpinned || out == PrefillRetireOutcome::None) { + slot.OnSuccessfulRetire(); + return true; + } + slot.quarantined = true; + return false; +} + +// Product-shaped Launch: arm rollback only after the first enqueue, then every +// later failure retires the streams that actually received work. +inline bool HostLaunch(PrefillDequantCacheHost& cache, PrefillPeerLife& slot, int device, + const void* key, int M, PrefillPeerFailAt fail, bool retire_ok = true, + bool restore_ok = true) { + if (!SlotReusable(slot.pending_M > 0, slot.rollback_armed, slot.quarantined) || + !slot.output_copy.CanReuseScratch()) + return false; + auto set = [restore_ok](int) { return restore_ok; }; + ComputeDevGuard guard(device, set); + slot.compute_restored = false; + + // First current-generation enqueue (compute-stream ev_c analogue). + slot.ArmRollback(); + slot.MarkComputeWork(); + if (fail == PrefillPeerFailAt::AfterFirstEnqueue) { + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + guard.RestoreOrThrow(); + slot.compute_restored = true; + return false; + } + slot.MarkExpertWork(); + + int pin = -1; + void* gu = nullptr; + void* dn = nullptr; + bool got = false; + { + std::lock_guard lk(cache.mu); + got = cache.Ensure(device, 4, 8) && cache.GetLocked(key, &gu, &dn, &pin); + if (!got && pin >= 0) { + slot.cache_pin = pin; + slot.cache_dev = device; + slot.fill_lease = true; + } + } + if (!got) { + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + guard.RestoreOrThrow(); + slot.compute_restored = true; + return false; + } + slot.cache_pin = pin; // persist immediately on acquisition + slot.cache_dev = device; + slot.fill_lease = false; + if (fail == PrefillPeerFailAt::AfterPinEnqueue || fail == PrefillPeerFailAt::RecordEvE) { + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + guard.RestoreOrThrow(); + slot.compute_restored = true; + return false; + } + slot.MarkThisGenEvent(); + if (fail == PrefillPeerFailAt::AfterRecord) { + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + guard.RestoreOrThrow(); + slot.compute_restored = true; + return false; + } + PublishThenRestoreOrThrow( + slot, M, [&] { guard.RestoreOrThrow(); }, + [&] { return HostRetireThenUnpin(cache, slot, retire_ok); }); + slot.compute_restored = true; + return true; +} + +inline bool HostFinish(PrefillDequantCacheHost& cache, PrefillPeerLife& slot, int M, + PrefillPeerFailAt fail, bool retire_ok = true, bool restore_ok = true, + int copy_stream = 0, int retire_stream = 0) { + if (slot.pending_M <= 0 || M > slot.pending_M) return false; + auto set = [restore_ok](int) { return restore_ok; }; + ComputeDevGuard guard(0, set); + if (fail == PrefillPeerFailAt::AfterWait) { + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + slot.pending_M = 0; + guard.RestoreOrThrow(); + slot.compute_restored = true; + return false; + } + slot.output_copy.Enqueue(copy_stream); + if (fail == PrefillPeerFailAt::AfterCopy) { + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + (void)slot.output_copy.Retire(copy_stream); + slot.pending_M = 0; + guard.RestoreOrThrow(); + slot.compute_restored = true; + return false; + } + if (!slot.output_copy.Retire(retire_stream)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + if (!HostRetireThenUnpin(cache, slot, retire_ok)) { + slot.compute_restored = true; + guard.RestoreOrThrow(); + return false; + } + slot.pending_M = 0; + guard.RestoreOrThrow(); + slot.compute_restored = true; + return true; +} + +// Same-device product seam: persist pin on acquire; GEMM readers run while +// pinned; unpin only after observed retire. Early UnpinLocked is a RED mutation. +struct SameDevSession { + PrefillDequantCacheHost* cache = nullptr; + SameDevLife* life = nullptr; + int pin = -1; + bool gemm_readers_ran = false; + bool unpinned_before_gemm = false; + + bool Acquire(const void* key, int device) { + if (!cache || !life || !life->CanEnter()) return false; + void* gu = nullptr; + void* dn = nullptr; + bool got = false; + { + std::lock_guard lk(cache->mu); + got = cache->Ensure(device, 4, 8) && cache->GetLocked(key, &gu, &dn, &pin); + if (!got && pin >= 0) life->PersistPin(pin, device); + } + if (!got) { + if (pin >= 0) { + std::lock_guard lk(cache->mu); + (void)cache->ReleaseObservedPinLocked(pin, /*fill_lease=*/true); + life->ClearPin(); + pin = -1; + } + return false; + } + life->PersistPin(pin, device); + return true; + } + + bool RunGemmReaders() { + if (!cache || pin < 0) return false; + if (cache->slots[pin].pins <= 0) { + unpinned_before_gemm = true; + return false; + } + gemm_readers_ran = true; + return true; + } + + bool Retire(bool sync_ok) { + if (!cache || !life) return false; + if (pin < 0) return true; + if (!sync_ok) { + life->Quarantine(pin, life->cache_dev); + return false; + } + { + std::lock_guard lk(cache->mu); + cache->UnpinLocked(pin); + } + life->ClearPin(); + pin = -1; + return true; + } +}; + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_gemma4_experts.hip b/src/vt/rocm/rocm_gemma4_experts.hip index e9facdf2e..7d0b51ce9 100644 --- a/src/vt/rocm/rocm_gemma4_experts.hip +++ b/src/vt/rocm/rocm_gemma4_experts.hip @@ -1,5 +1,6 @@ // ROCm upload of Gemma-4 MoE expert stacks (BF16 fused or FP8→BF16 dequant). #include +#include "vt/rocm/rocm_gemma4_prefill_dequant_cache.h" // Packs GPU0 first (compute device), then GPU1. FP8 streams per-expert to // avoid ~30G host OOM from permanent BF16 caches / full-layer host buffers. #include @@ -9,6 +10,8 @@ #include #include #include +#include +#include #include #include "vllm/model_executor/models/gemma4.h" @@ -645,6 +648,166 @@ bool RunGemma4Fp8TopKIndexedOnExpertDevice(vt::Queue& compute_q, int expert_dev, // Prefill: run ExpertGeGLU for M rows on expert_dev (weights resident there). // Peer only activations M×H bf16 — not multi-GB FP8 expert packs. +// #839: wrapper → Launch/Finish + pin until host-observed retirement. + +struct HipPrefillCacheHooks { + vt::Queue* q = nullptr; + const void* fp8_gu = nullptr; + const void* s_gu = nullptr; + const void* fp8_dn = nullptr; + const void* s_dn = nullptr; + int I = 0, H = 0; + struct Event { + hipEvent_t h = nullptr; + }; + bool SetDevice(int d) { return hipSetDevice(d) == hipSuccess; } + void* Malloc(size_t n) { + void* p = nullptr; + return hipMalloc(&p, n) == hipSuccess ? p : nullptr; + } + void Free(void* p) { + if (p) (void)hipFree(p); + } + void DestroyEvent(Event& e) { + if (e.h) { + (void)hipEventDestroy(e.h); + e.h = nullptr; + } + } + bool Fill(void* gu, void* dn) { + if (!q || !fp8_gu || !s_gu || !fp8_dn || !s_dn || I <= 0 || H <= 0 || !gu || !dn) return false; + vt::DequantFp8ChannelBf16(*q, gu, fp8_gu, s_gu, 2 * I, H); + vt::DequantFp8ChannelBf16(*q, dn, fp8_dn, s_dn, H, I); + return true; + } + bool RecordReady(Event& e) { + if (!q) return false; + if (!e.h) { + if (hipEventCreateWithFlags(&e.h, hipEventDisableTiming) != hipSuccess) return false; + } + return hipEventRecord(e.h, static_cast(q->handle)) == hipSuccess; + } + bool WaitReady(Event& e) { + if (!q || !e.h) return false; + return hipStreamWaitEvent(static_cast(q->handle), e.h, 0) == hipSuccess; + } +}; + +using PrefillDequantCache = vt::rocm::PrefillDequantCacheT; + +inline bool BindPrefillCacheFill(PrefillDequantCache& cache, vt::Queue& q, const void* fp8_gu, + const void* s_gu, const void* fp8_dn, const void* s_dn, int I, + int H) { + cache.hooks.q = &q; + cache.hooks.fp8_gu = fp8_gu; + cache.hooks.s_gu = s_gu; + cache.hooks.fp8_dn = fp8_dn; + cache.hooks.s_dn = s_dn; + cache.hooks.I = I; + cache.hooks.H = H; + return true; +} + +struct DequantCacheSlot { + PrefillDequantCache cache; +}; + +DequantCacheSlot& DequantCacheSlotFor(int device) { + static DequantCacheSlot slots[8]; + if (device < 0 || device >= 8) { + static DequantCacheSlot fallback; + return fallback; + } + return slots[device]; +} + +struct PeerSlot { + int cdev = -1, edev = -1; + int I = 0, H = 0, Mcap = 0; + int pending_M = 0; + vt::Queue eq{}; + bool eq_live = false; + void* x = nullptr; + void* y = nullptr; + void* gu_act = nullptr; + void* act = nullptr; + hipEvent_t ev_c = nullptr, ev_e = nullptr; + int cache_pin = -1; + int cache_dev = -1; + bool ev_e_recorded = false; + bool this_gen_ev_e = false; + bool work_on_compute = false; + bool work_on_expert = false; + bool rollback_armed = false; + bool quarantined = false; + bool fill_lease = false; +}; + +struct PeerPipe { + PeerSlot s[1]; // slot 0 only; overlap slot-1 is a later row +}; + +inline PeerPipe& PeerPipeTls() { + static thread_local PeerPipe pipe; + return pipe; +} + +vt::rocm::PrefillPeerLife PeerLifeView(const PeerSlot& tls) { return vt::rocm::LifeFromSlot(tls); } + +void RestoreComputeOrThrow(int compute_dev) { + if (hipSetDevice(compute_dev) != hipSuccess) throw vt::rocm::RestoreFailed{}; +} + +// Unpin only after a successful event/stream retirement. Failure quarantines the pin. +// ChoosePrefillRetire forbids syncing a leftover prior ev_e during current-gen rollback. +bool RetirePinThenUnpin(PeerSlot& tls, hipStream_t expert_stream, hipStream_t compute_stream) { + const auto target = vt::rocm::ChoosePrefillRetire(PeerLifeView(tls)); + bool retire_ok = true; + if (target == vt::rocm::PrefillRetireTarget::RecordedEvent && tls.ev_e) { + retire_ok = hipEventSynchronize(tls.ev_e) == hipSuccess; + } else if (target == vt::rocm::PrefillRetireTarget::ExpertStream && expert_stream) { + retire_ok = hipStreamSynchronize(expert_stream) == hipSuccess; + } else if (target == vt::rocm::PrefillRetireTarget::ComputeStream && compute_stream) { + retire_ok = hipStreamSynchronize(compute_stream) == hipSuccess; + } else if (target == vt::rocm::PrefillRetireTarget::ComputeStream && expert_stream) { + // no compute handle (should not happen after first enqueue) — do not fake-succeed + retire_ok = false; + } + if (!retire_ok) { + tls.quarantined = true; + return false; + } + if (tls.cache_pin >= 0 && tls.cache_dev >= 0) { + auto& wslot = DequantCacheSlotFor(tls.cache_dev); + std::lock_guard lk(wslot.cache.mu); + if (!wslot.cache.ReleaseObservedPinLocked(tls.cache_pin, tls.fill_lease)) { + tls.quarantined = true; + return false; + } + } + tls.cache_pin = -1; + tls.cache_dev = -1; + tls.rollback_armed = false; + tls.ev_e_recorded = false; + tls.this_gen_ev_e = false; + tls.work_on_compute = false; + tls.work_on_expert = false; + tls.fill_lease = false; + return true; +} + +void FailLaunchRestore(PeerSlot& tls, hipStream_t est, hipStream_t cst, int compute_dev) { + if (!RetirePinThenUnpin(tls, est, cst)) tls.quarantined = true; + RestoreComputeOrThrow(compute_dev); +} + +bool LaunchGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, int expert_dev, + const void* x_compute, const void* fp8_gu, + const void* s_gu, const void* fp8_dn, const void* s_dn, + int M, int I, int H); +bool FinishGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, void* y_compute, int M, + int H); + bool RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice(vt::Queue& compute_q, int expert_dev, void* y_compute, const void* x_compute, const void* fp8_gu, const void* s_gu, @@ -655,101 +818,145 @@ bool RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice(vt::Queue& compute_q, int expe !fp8_dn || !s_dn || M <= 0 || I <= 0 || H <= 0) return false; - // Same device: caller should use local ExpertGeGLU; still support for safety. if (expert_dev == compute_dev) { vt::Device dev = compute_q.device; - // Dequant sticky into TLS then GEMM on compute_q. struct SameTls { int dev = -1; int I = 0, H = 0, Mcap = 0; - const void* sticky = nullptr; - void* gu_w = nullptr; - void* dn_w = nullptr; void* gu_act = nullptr; void* act = nullptr; + vt::rocm::SameDevLife life{}; }; - static thread_local SameTls tls; - const size_t gu_w_b = static_cast(2 * I) * static_cast(H) * 2; - const size_t dn_w_b = static_cast(H) * static_cast(I) * 2; - const size_t gu_act_b = static_cast(M) * static_cast(2 * I) * 2; - const size_t act_b = static_cast(M) * static_cast(I) * 2; - if (tls.dev != compute_dev || tls.I != I || tls.H != H || tls.Mcap < M || !tls.gu_w) { + static thread_local SameTls tls_slots[2]; + SameTls& tls = tls_slots[(compute_dev == 1) ? 1 : 0]; + if (!tls.life.CanEnter()) return false; + if ((tls.dev != compute_dev || tls.I != I || tls.H != H || tls.Mcap < M || !tls.gu_act) && + !tls.life.CanReconfigure()) + return false; + if (tls.dev != compute_dev || tls.I != I || tls.H != H || tls.Mcap < M || !tls.gu_act) { auto freep = [&](void*& p) { if (p) { (void)hipFree(p); p = nullptr; } }; - freep(tls.gu_w); - freep(tls.dn_w); freep(tls.gu_act); freep(tls.act); Check(hipSetDevice(compute_dev), "same alloc"); - if (hipMalloc(&tls.gu_w, gu_w_b) != hipSuccess) return false; - if (hipMalloc(&tls.dn_w, dn_w_b) != hipSuccess) return false; - // Cap act buffers to at least M; grow only. size_t gu_act_cap = static_cast(std::max(M, 64)) * static_cast(2 * I) * 2; size_t act_cap = static_cast(std::max(M, 64)) * static_cast(I) * 2; - if (hipMalloc(&tls.gu_act, gu_act_cap) != hipSuccess) return false; - if (hipMalloc(&tls.act, act_cap) != hipSuccess) return false; + void* ngu = nullptr; + void* nact = nullptr; + if (hipMalloc(&ngu, gu_act_cap) != hipSuccess) return false; + if (hipMalloc(&nact, act_cap) != hipSuccess) { + (void)hipFree(ngu); + return false; + } + tls.gu_act = ngu; + tls.act = nact; tls.dev = compute_dev; tls.I = I; tls.H = H; tls.Mcap = std::max(M, 64); - tls.sticky = nullptr; - } - if (tls.sticky != fp8_gu) { - vt::DequantFp8ChannelBf16(compute_q, tls.gu_w, fp8_gu, s_gu, 2 * I, H); - vt::DequantFp8ChannelBf16(compute_q, tls.dn_w, fp8_dn, s_dn, H, I); - tls.sticky = fp8_gu; - } - vt::Tensor x = vt::Tensor::Contiguous(const_cast(x_compute), vt::DType::kBF16, dev, - {M, H}); - vt::Tensor gu_w = vt::Tensor::Contiguous(tls.gu_w, vt::DType::kBF16, dev, {2 * I, H}); - vt::Tensor gu_act = vt::Tensor::Contiguous(tls.gu_act, vt::DType::kBF16, dev, {M, 2 * I}); - vt::Tensor act = vt::Tensor::Contiguous(tls.act, vt::DType::kBF16, dev, {M, I}); - vt::MatmulBT(compute_q, gu_act, x, gu_w); - vt::GeluAndMul(compute_q, act, gu_act); - vt::MatmulBTAlphaBeta(compute_q, y_compute, act.data, tls.dn_w, M, H, I, 1.f, 0.f, - vt::DType::kBF16); - return true; - } - - struct Tls { - int cdev = -1, edev = -1; - int I = 0, H = 0, Mcap = 0; - const void* sticky = nullptr; - vt::Queue eq{}; - bool eq_live = false; - void* x = nullptr; - void* y = nullptr; + } + auto& wslot = DequantCacheSlotFor(compute_dev); void* gu_w = nullptr; void* dn_w = nullptr; - void* gu_act = nullptr; - void* act = nullptr; - hipEvent_t ev_c = nullptr, ev_e = nullptr; - }; - static thread_local Tls tls; - const size_t row_b = static_cast(H) * 2; - const size_t xy_need = static_cast(M) * row_b; - const size_t gu_w_b = static_cast(2 * I) * static_cast(H) * 2; - const size_t dn_w_b = static_cast(H) * static_cast(I) * 2; - hipStream_t cst = static_cast(compute_q.handle); - vt::Device edev{vt::DeviceType::kROCM, expert_dev}; + int pin = -1; + hipStream_t cst = static_cast(compute_q.handle); + { + std::lock_guard lk(wslot.cache.mu); + if (!wslot.cache.Ensure(compute_dev, I, H)) return false; + BindPrefillCacheFill(wslot.cache, compute_q, fp8_gu, s_gu, fp8_dn, s_dn, I, H); + if (!wslot.cache.GetLocked(fp8_gu, &gu_w, &dn_w, &pin)) { + if (pin >= 0) tls.life.PersistPin(pin, compute_dev); + } + } + if (pin >= 0 && (gu_w == nullptr || dn_w == nullptr)) { + const bool ok = hipStreamSynchronize(cst) == hipSuccess; + if (!ok) { + tls.life.Quarantine(pin, compute_dev); + return false; + } + { + std::lock_guard lk(wslot.cache.mu); + if (!wslot.cache.RetireFillLocked(pin, true)) { + tls.life.Quarantine(pin, compute_dev); + return false; + } + } + tls.life.ClearPin(); + return false; + } + if (gu_w == nullptr || dn_w == nullptr) return false; + tls.life.PersistPin(pin, compute_dev); + try { + vt::Tensor x = vt::Tensor::Contiguous(const_cast(x_compute), vt::DType::kBF16, dev, {M, H}); + vt::Tensor gu_wt = vt::Tensor::Contiguous(gu_w, vt::DType::kBF16, dev, {2 * I, H}); + vt::Tensor gu_act = vt::Tensor::Contiguous(tls.gu_act, vt::DType::kBF16, dev, {M, 2 * I}); + vt::Tensor act = vt::Tensor::Contiguous(tls.act, vt::DType::kBF16, dev, {M, I}); + vt::MatmulBT(compute_q, gu_act, x, gu_wt); + vt::GeluAndMul(compute_q, act, gu_act); + vt::MatmulBTAlphaBeta(compute_q, y_compute, act.data, dn_w, M, H, I, 1.f, 0.f, vt::DType::kBF16); + if (hipStreamSynchronize(cst) != hipSuccess) { + tls.life.Quarantine(pin, compute_dev); + RestoreComputeOrThrow(compute_dev); + return false; + } + { + std::lock_guard lk(wslot.cache.mu); + wslot.cache.UnpinLocked(pin); + } + tls.life.ClearPin(); + RestoreComputeOrThrow(compute_dev); + return true; + } catch (...) { + if (hipStreamSynchronize(cst) != hipSuccess) tls.life.Quarantine(pin, compute_dev); + else { + std::lock_guard lk(wslot.cache.mu); + wslot.cache.UnpinLocked(pin); + tls.life.ClearPin(); + } + RestoreComputeOrThrow(compute_dev); + throw; + } + } + + if (!LaunchGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0, expert_dev, x_compute, fp8_gu, + s_gu, fp8_dn, s_dn, M, I, H)) + return false; + return FinishGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0, y_compute, M, H); +} +bool LaunchGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, int expert_dev, + const void* x_compute, const void* fp8_gu, + const void* s_gu, const void* fp8_dn, const void* s_dn, + int M, int I, int H) { + try { + if (slot != 0 || !x_compute || !fp8_gu || !s_gu || !fp8_dn || !s_dn || M <= 0 || + I <= 0 || H <= 0 || expert_dev < 0) + return false; + const int compute_dev = compute_q.device.index; + if (compute_dev < 0 || expert_dev == compute_dev) return false; + + PeerSlot& tls = PeerPipeTls().s[slot]; auto freep = [](int dev, void*& p) { if (!p) return; (void)hipSetDevice(dev); (void)hipFree(p); p = nullptr; }; + const size_t row_b = static_cast(H) * 2; + const size_t xy_need = static_cast(M) * row_b; + vt::Device edev{vt::DeviceType::kROCM, expert_dev}; + hipStream_t cst = static_cast(compute_q.handle); + + if (!vt::rocm::SlotReusable(tls.pending_M > 0, tls.rollback_armed, tls.quarantined)) return false; if (tls.edev != expert_dev || tls.cdev != compute_dev || tls.I != I || tls.H != H || tls.Mcap < M || !tls.x) { freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.x); freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.y); - freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.gu_w); - freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.dn_w); freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.gu_act); freep(tls.edev >= 0 ? tls.edev : expert_dev, tls.act); if (tls.eq_live) { @@ -765,82 +972,235 @@ bool RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice(vt::Queue& compute_q, int expe tls.ev_e = nullptr; } const int Mcap = std::max(M, 64); - Check(hipSetDevice(expert_dev), "peer-act expert alloc"); - if (hipMalloc(&tls.x, static_cast(Mcap) * row_b) != hipSuccess) return false; - if (hipMalloc(&tls.y, static_cast(Mcap) * row_b) != hipSuccess) return false; - if (hipMalloc(&tls.gu_w, gu_w_b) != hipSuccess) return false; - if (hipMalloc(&tls.dn_w, dn_w_b) != hipSuccess) return false; - if (hipMalloc(&tls.gu_act, static_cast(Mcap) * static_cast(2 * I) * 2) != - hipSuccess) - return false; - if (hipMalloc(&tls.act, static_cast(Mcap) * static_cast(I) * 2) != hipSuccess) - return false; - tls.eq = vt::CreateQueue(edev); + void* nx = nullptr; + void* ny = nullptr; + void* ngu = nullptr; + void* nact = nullptr; + hipEvent_t nev_c = nullptr; + hipEvent_t nev_e = nullptr; + vt::Queue nq{}; + bool nq_live = false; + auto drop = [&]() { + if (nx) (void)hipFree(nx); + if (ny) (void)hipFree(ny); + if (ngu) (void)hipFree(ngu); + if (nact) (void)hipFree(nact); + if (nev_c) (void)hipEventDestroy(nev_c); + if (nev_e) (void)hipEventDestroy(nev_e); + if (nq_live) vt::DestroyQueue(nq); + }; + try { + Check(hipSetDevice(expert_dev), "peer-pipe expert alloc"); + if (hipMalloc(&nx, static_cast(Mcap) * row_b) != hipSuccess || + hipMalloc(&ny, static_cast(Mcap) * row_b) != hipSuccess || + hipMalloc(&ngu, static_cast(Mcap) * static_cast(2 * I) * 2) != hipSuccess || + hipMalloc(&nact, static_cast(Mcap) * static_cast(I) * 2) != hipSuccess) { + drop(); + RestoreComputeOrThrow(compute_dev); + return false; + } + nq = vt::CreateQueue(edev); + nq_live = true; + Check(hipSetDevice(compute_dev), "peer-pipe ev_c"); + Check(hipEventCreateWithFlags(&nev_c, hipEventDisableTiming), "ev_c"); + Check(hipSetDevice(expert_dev), "peer-pipe ev_e"); + Check(hipEventCreateWithFlags(&nev_e, hipEventDisableTiming), "ev_e"); + Check(hipSetDevice(compute_dev), "p2p c"); + EnablePeerAccessQuiet(expert_dev); + Check(hipSetDevice(expert_dev), "p2p e"); + EnablePeerAccessQuiet(compute_dev); + } catch (...) { + drop(); + RestoreComputeOrThrow(compute_dev); + throw; + } + tls.x = nx; + tls.y = ny; + tls.gu_act = ngu; + tls.act = nact; + tls.eq = nq; tls.eq_live = true; - Check(hipSetDevice(compute_dev), "peer-act ev_c"); - Check(hipEventCreateWithFlags(&tls.ev_c, hipEventDisableTiming), "ev_c"); - Check(hipSetDevice(expert_dev), "peer-act ev_e"); - Check(hipEventCreateWithFlags(&tls.ev_e, hipEventDisableTiming), "ev_e"); - Check(hipSetDevice(compute_dev), "p2p c"); - EnablePeerAccessQuiet(expert_dev); - Check(hipSetDevice(expert_dev), "p2p e"); - EnablePeerAccessQuiet(compute_dev); + tls.ev_c = nev_c; + tls.ev_e = nev_e; tls.edev = expert_dev; tls.cdev = compute_dev; tls.I = I; tls.H = H; tls.Mcap = Mcap; - tls.sticky = nullptr; + tls.cache_pin = -1; + tls.cache_dev = -1; + tls.ev_e_recorded = false; + tls.this_gen_ev_e = false; + tls.work_on_compute = false; + tls.work_on_expert = false; + tls.rollback_armed = false; } hipStream_t est = static_cast(tls.eq.handle); - if (hipEventRecord(tls.ev_c, cst) != hipSuccess) return false; - if (hipStreamWaitEvent(est, tls.ev_c, 0) != hipSuccess) return false; + if (hipEventRecord(tls.ev_c, cst) != hipSuccess) { + RestoreComputeOrThrow(compute_dev); return false; + } + // First successful current-gen enqueue is on the compute stream. + tls.rollback_armed = true; + tls.work_on_compute = true; + tls.this_gen_ev_e = false; + if (hipStreamWaitEvent(est, tls.ev_c, 0) != hipSuccess) { + FailLaunchRestore(tls, est, cst, compute_dev); return false; + } + tls.work_on_expert = true; bool x_ok = hipMemcpyPeerAsync(tls.x, expert_dev, const_cast(x_compute), compute_dev, xy_need, est) == hipSuccess; if (!x_ok) { std::vector tmp(xy_need); Check(hipSetDevice(compute_dev), "x d2h"); - if (hipStreamSynchronize(cst) != hipSuccess) return false; - if (hipMemcpy(tmp.data(), x_compute, xy_need, hipMemcpyDeviceToHost) != hipSuccess) - return false; + if (hipStreamSynchronize(cst) != hipSuccess) { + FailLaunchRestore(tls, est, cst, compute_dev); return false; + } + if (hipMemcpy(tmp.data(), x_compute, xy_need, hipMemcpyDeviceToHost) != hipSuccess) { + FailLaunchRestore(tls, est, cst, compute_dev); return false; + } Check(hipSetDevice(expert_dev), "x h2d"); - if (hipMemcpyAsync(tls.x, tmp.data(), xy_need, hipMemcpyHostToDevice, est) != hipSuccess) - return false; + if (hipMemcpyAsync(tls.x, tmp.data(), xy_need, hipMemcpyHostToDevice, est) != hipSuccess) { + FailLaunchRestore(tls, est, cst, compute_dev); return false; + } } - if (tls.sticky != fp8_gu) { - Check(hipSetDevice(expert_dev), "dequant setdev"); - vt::DequantFp8ChannelBf16(tls.eq, tls.gu_w, fp8_gu, s_gu, 2 * I, H); - vt::DequantFp8ChannelBf16(tls.eq, tls.dn_w, fp8_dn, s_dn, H, I); - tls.sticky = fp8_gu; + auto& wslot = DequantCacheSlotFor(expert_dev); + void* gu_w = nullptr; + void* dn_w = nullptr; + int pin = -1; + bool got = false; + { + std::lock_guard lk(wslot.cache.mu); + if (!wslot.cache.Ensure(expert_dev, I, H)) { + got = false; + } else { + Check(hipSetDevice(expert_dev), "dequant setdev"); + BindPrefillCacheFill(wslot.cache, tls.eq, fp8_gu, s_gu, fp8_dn, s_dn, I, H); + got = wslot.cache.GetLocked(fp8_gu, &gu_w, &dn_w, &pin); + if (!got && pin >= 0) { + tls.cache_pin = pin; + tls.cache_dev = expert_dev; + tls.fill_lease = true; + } + } + } + if (!got) { + FailLaunchRestore(tls, est, cst, compute_dev); return false; } + tls.cache_pin = pin; + tls.cache_dev = expert_dev; + tls.fill_lease = false; + tls.ev_e_recorded = false; - Check(hipSetDevice(expert_dev), "geglu setdev"); vt::Tensor x = vt::Tensor::Contiguous(tls.x, vt::DType::kBF16, edev, {M, H}); - vt::Tensor gu_w = vt::Tensor::Contiguous(tls.gu_w, vt::DType::kBF16, edev, {2 * I, H}); + vt::Tensor gu_wt = vt::Tensor::Contiguous(gu_w, vt::DType::kBF16, edev, {2 * I, H}); vt::Tensor gu_act = vt::Tensor::Contiguous(tls.gu_act, vt::DType::kBF16, edev, {M, 2 * I}); vt::Tensor act = vt::Tensor::Contiguous(tls.act, vt::DType::kBF16, edev, {M, I}); - vt::MatmulBT(tls.eq, gu_act, x, gu_w); + vt::MatmulBT(tls.eq, gu_act, x, gu_wt); vt::GeluAndMul(tls.eq, act, gu_act); - vt::MatmulBTAlphaBeta(tls.eq, tls.y, act.data, tls.dn_w, M, H, I, 1.f, 0.f, vt::DType::kBF16); + vt::MatmulBTAlphaBeta(tls.eq, tls.y, act.data, dn_w, M, H, I, 1.f, 0.f, vt::DType::kBF16); - if (hipEventRecord(tls.ev_e, est) != hipSuccess) return false; - if (hipStreamWaitEvent(cst, tls.ev_e, 0) != hipSuccess) return false; + if (hipEventRecord(tls.ev_e, est) != hipSuccess) { + FailLaunchRestore(tls, est, cst, compute_dev); return false; + } + tls.this_gen_ev_e = true; + tls.ev_e_recorded = true; + vt::rocm::PublishThenRestoreOrThrow( + tls, M, [&] { RestoreComputeOrThrow(compute_dev); }, + [&] { return RetirePinThenUnpin(tls, est, cst); }); + return true; + } catch (const vt::rocm::RestoreFailed&) { + PeerSlot& tls = PeerPipeTls().s[slot]; + hipStream_t est = tls.eq_live ? static_cast(tls.eq.handle) : nullptr; + hipStream_t cst = static_cast(compute_q.handle); + if (tls.cache_pin >= 0 || tls.rollback_armed || tls.pending_M > 0) + (void)RetirePinThenUnpin(tls, est, cst); + tls.pending_M = 0; + throw; + } catch (...) { + const int compute_dev = compute_q.device.index; + PeerSlot& tls = PeerPipeTls().s[slot]; + hipStream_t est = tls.eq_live ? static_cast(tls.eq.handle) : nullptr; + hipStream_t cst = static_cast(compute_q.handle); + if (tls.rollback_armed || tls.cache_pin >= 0) (void)RetirePinThenUnpin(tls, est, cst); + RestoreComputeOrThrow(compute_dev); + throw; + } +} + +bool FinishGemma4Fp8ExpertGeGLUPrefillPeer(vt::Queue& compute_q, int slot, void* y_compute, int M, + int H) { + try { + if (slot != 0 || !y_compute || M <= 0 || H <= 0) return false; + PeerSlot& tls = PeerPipeTls().s[slot]; + if (tls.pending_M <= 0 || !tls.y || !tls.ev_e) return false; + if (M > tls.pending_M) return false; + const int compute_dev = compute_q.device.index; + const int expert_dev = tls.edev; + const size_t xy_need = static_cast(M) * static_cast(H) * 2; + hipStream_t cst = static_cast(compute_q.handle); + hipStream_t est = tls.eq_live ? static_cast(tls.eq.handle) : nullptr; + + if (hipStreamWaitEvent(cst, tls.ev_e, 0) != hipSuccess) { + if (!RetirePinThenUnpin(tls, est, cst)) tls.quarantined = true; + else tls.pending_M = 0; + RestoreComputeOrThrow(compute_dev); return false; + } bool y_ok = hipMemcpyPeerAsync(y_compute, compute_dev, tls.y, expert_dev, xy_need, cst) == hipSuccess; if (!y_ok) { std::vector tmp(xy_need); Check(hipSetDevice(expert_dev), "y d2h"); - if (hipStreamSynchronize(est) != hipSuccess) return false; - if (hipMemcpy(tmp.data(), tls.y, xy_need, hipMemcpyDeviceToHost) != hipSuccess) return false; + if (est && hipStreamSynchronize(est) != hipSuccess) { + tls.quarantined = true; + RestoreComputeOrThrow(compute_dev); return false; + } + if (hipMemcpy(tmp.data(), tls.y, xy_need, hipMemcpyDeviceToHost) != hipSuccess) { + if (!RetirePinThenUnpin(tls, est, cst)) tls.quarantined = true; + else tls.pending_M = 0; + RestoreComputeOrThrow(compute_dev); return false; + } Check(hipSetDevice(compute_dev), "y h2d"); - if (hipMemcpyAsync(y_compute, tmp.data(), xy_need, hipMemcpyHostToDevice, cst) != hipSuccess) - return false; + if (hipMemcpyAsync(y_compute, tmp.data(), xy_need, hipMemcpyHostToDevice, cst) != hipSuccess) { + if (!RetirePinThenUnpin(tls, est, cst)) tls.quarantined = true; + else tls.pending_M = 0; + RestoreComputeOrThrow(compute_dev); return false; + } + } + if (tls.ev_e && hipEventSynchronize(tls.ev_e) != hipSuccess) { + tls.quarantined = true; + RestoreComputeOrThrow(compute_dev); return false; + } + tls.ev_e_recorded = true; + // Output copy is on cst. Host-observe that stream before pending/reuse. + if (hipStreamSynchronize(cst) != hipSuccess) { + tls.quarantined = true; + RestoreComputeOrThrow(compute_dev); return false; + } + if (!RetirePinThenUnpin(tls, est, cst)) { + tls.quarantined = true; + RestoreComputeOrThrow(compute_dev); return false; + } + tls.pending_M = 0; + RestoreComputeOrThrow(compute_dev); return true; + } catch (const vt::rocm::RestoreFailed&) { + PeerSlot& tls = PeerPipeTls().s[slot]; + hipStream_t est = tls.eq_live ? static_cast(tls.eq.handle) : nullptr; + hipStream_t cst = static_cast(compute_q.handle); + if (tls.cache_pin >= 0 || tls.rollback_armed) (void)RetirePinThenUnpin(tls, est, cst); + tls.pending_M = 0; + throw; + } catch (...) { + const int compute_dev = compute_q.device.index; + PeerSlot& tls = PeerPipeTls().s[slot]; + hipStream_t est = tls.eq_live ? static_cast(tls.eq.handle) : nullptr; + hipStream_t cst = static_cast(compute_q.handle); + if (tls.cache_pin >= 0 || tls.rollback_armed) (void)RetirePinThenUnpin(tls, est, cst); + RestoreComputeOrThrow(compute_dev); + throw; } - if (hipSetDevice(compute_dev) != hipSuccess) return false; - return true; } void PinGemma4Fp8ExpertHostCache(const Gemma4Fp8ExpertMats& ex) { diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index b8c052162..a1c43a8ca 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1641,6 +1641,10 @@ target_link_libraries(test_parallel_layout PRIVATE Threads::Threads) vllm_cpp_add_test(test_backend_multidevice vt/test_backend_multidevice.cpp) vllm_cpp_add_test(test_arena vt/test_arena.cpp) vllm_cpp_add_test(test_ops_matmul vt/test_ops_matmul.cpp) +# #839 host lifetime + source invariants (no GPU). +vllm_cpp_add_test(test_ops_gemma4_prefill_peer vt/test_ops_gemma4_prefill_peer.cpp) +target_compile_definitions(test_ops_gemma4_prefill_peer PRIVATE + VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") # D1 (fold plan Tier D1): the merged-QKV == separate-q/k/v A/B — one MatmulBT over # the merged [Nq+Nk+Nv,H] owner + a contiguous QkvSplit is BYTE-IDENTICAL to three # separate MatmulBT (bit-exact GEMM math, wider N; RED-first). CPU-only. diff --git a/tests/vt/test_ops_gemma4_prefill_peer.cpp b/tests/vt/test_ops_gemma4_prefill_peer.cpp new file mode 100644 index 000000000..2c57dcf81 --- /dev/null +++ b/tests/vt/test_ops_gemma4_prefill_peer.cpp @@ -0,0 +1,686 @@ +// #839 host lifetime + product-policy mutations for prefill peer. +#include +#include +#include +#include +#include +#include + +#include + +#include "vt/rocm/rocm_gemma4_prefill_dequant_cache.h" + +#ifndef VLLM_CPP_SOURCE_DIR +#define VLLM_CPP_SOURCE_DIR "." +#endif + +namespace { + +std::string ReadHip() { + const std::string path = std::string(VLLM_CPP_SOURCE_DIR) + "/src/vt/rocm/rocm_gemma4_experts.hip"; + std::ifstream in(path); + REQUIRE(in.good()); + return std::string((std::istreambuf_iterator(in)), std::istreambuf_iterator()); +} + +std::string ReadHeader() { + const std::string path = + std::string(VLLM_CPP_SOURCE_DIR) + "/include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h"; + std::ifstream in(path); + REQUIRE(in.good()); + return std::string((std::istreambuf_iterator(in)), std::istreambuf_iterator()); +} + +std::string ExtractFinish(const std::string& hip) { + const auto a = hip.rfind("bool FinishGemma4Fp8ExpertGeGLUPrefillPeer"); + REQUIRE(a != std::string::npos); + const auto b = hip.find("void PinGemma4Fp8ExpertHostCache", a); + REQUIRE(b != std::string::npos); + return hip.substr(a, b - a); +} + +std::string ExtractRestoreComputeOrThrow(const std::string& hip) { + const auto a = hip.find("void RestoreComputeOrThrow(int compute_dev)"); + REQUIRE(a != std::string::npos); + const auto b = hip.find("\nbool RetirePinThenUnpin", a); + REQUIRE(b != std::string::npos); + return hip.substr(a, b - a); +} + +// The HIP half of the two compile gates below needs a HIP compiler. This file +// used to name one absolute path from the contributor's box +// ("/opt/rocm-7.2.4/core-7.14/bin/hipcc") and CHECK that it returned 0, so on +// every machine without that exact path the gate FAILED instead of reporting +// that it had not run -- the whole build-test-cpu and both sanitize-cpu legs of +// #1047, green on main. Resolve the toolchain the way +// tests/vt/test_ops_getblas_product.cpp:15-21 resolves its precondition, and say +// out loud when it is absent. +// +// Absent hipcc does NOT exit 77 here, unlike getblas: that file is one HIP gate +// and nothing else, while this one carries 21 host-lifetime cases that are the +// only gate CI has on this change. Exiting would take those with it. The g++ leg +// of both compile gates still runs and still asserts, so neither case reports a +// zero-assertion pass; what is lost off ROCm is the HIP-header compile, which no +// CI runner here can perform. See `## Owed` in the spec. +const char* ResolveHipcc() { + static const std::string resolved = [] { + std::vector candidates; + for (const char* var : {"VLLM_CPP_HIPCC", "HIPCC"}) { + if (const char* v = std::getenv(var); v != nullptr && v[0] != '\0') candidates.emplace_back(v); + } + if (const char* rp = std::getenv("ROCM_PATH"); rp != nullptr && rp[0] != '\0') { + candidates.emplace_back(std::string(rp) + "/bin/hipcc"); + } + candidates.emplace_back("hipcc"); // PATH + candidates.emplace_back("/opt/rocm/bin/hipcc"); + for (const auto& c : candidates) { + const std::string probe = c + " --version >/dev/null 2>&1"; + if (std::system(probe.c_str()) == 0) return c; + } + return std::string{}; + }(); + return resolved.empty() ? nullptr : resolved.c_str(); +} + +// Loud, and on stderr, so a reader of a CI log can tell a gate that ran from one +// that could not. Never silent: a skip nobody can see is the failure mode this +// repository's exit-77 convention exists to prevent. +void ReportHipccNotRun(const char* what) { + std::fprintf(stderr, + "\n*** HIP COMPILE GATE NOT RUN — no hipcc resolved, this is NOT a pass ***\n" + "%s\n" + "Set VLLM_CPP_HIPCC=/path/to/hipcc (or put hipcc on PATH, or set ROCM_PATH) to run it.\n", + what); +} + +struct RestoreGateRc { + int compile = 127; + int run = 127; +}; + +RestoreGateRc CompileAndRunProductRestore(const std::string& restore_fn, const std::string& compiler) { + const std::string path = "/tmp/vllm_prefill_peer_restore_gate.cpp"; + const std::string bin = "/tmp/vllm_prefill_peer_restore_gate.bin"; + std::ofstream out(path); + if (!out.good()) return {126, 126}; + out << R"GATE( +#include +namespace vt { +namespace rocm { +struct RestoreFailed : std::runtime_error { + RestoreFailed() : std::runtime_error("restore failed") {} +}; +} +} +using hipError_t = int; +constexpr hipError_t hipSuccess = 0; +static hipError_t g_set_rc = 1; +hipError_t hipSetDevice(int) { return g_set_rc; } +)GATE"; + out << restore_fn << '\n'; + out << R"GATE( +int probe() { + try { + RestoreComputeOrThrow(3); + return 0; + } catch (const vt::rocm::RestoreFailed&) { + return 1; + } catch (...) { + return 2; + } +} +int main() { + g_set_rc = 1; + if (probe() != 1) return 10; + g_set_rc = 0; + if (probe() != 0) return 11; + return 0; +} +)GATE"; + out.close(); + const std::string cmd = std::string("HIP_VISIBLE_DEVICES= ") + compiler + " -std=c++17 " + path + + " -o " + bin + + " >/tmp/vllm_prefill_peer_restore_gate.out 2>/tmp/vllm_prefill_peer_restore_gate.err"; + const int crc = std::system(cmd.c_str()); + if (crc != 0) return {crc, 127}; + return {0, std::system(bin.c_str())}; +} + +bool ProductFinishRetiresCstBeforeReuse(const std::string& finish) { + const auto sync = finish.find("hipStreamSynchronize(cst)"); + const auto success = + finish.find("tls.pending_M = 0;\n RestoreComputeOrThrow(compute_dev); return true;"); + return sync != std::string::npos && success != std::string::npos && sync < success; +} + +int CompileFinishCatchTu(const std::string& finish, const std::string& compiler) { + const auto c1 = finish.find("} catch (const vt::rocm::RestoreFailed&)"); + if (c1 == std::string::npos) return 127; + const std::string catches = finish.substr(c1); // includes function closer + const std::string path = "/tmp/vllm_prefill_peer_finish_catch_gate.cpp"; + std::ofstream out(path); + if (!out.good()) return 126; + out << R"GATE( +#include +namespace vt { +struct Queue { + struct { + int index = 0; + } device; + void* handle = nullptr; +}; +namespace rocm { +struct RestoreFailed : std::exception {}; +} +} +struct PeerSlot { + int cache_pin = -1; + bool rollback_armed = false; + int pending_M = 0; + bool eq_live = false; + struct { + void* handle = nullptr; + } eq; +}; +struct PeerTls { + PeerSlot s[1]; +}; +PeerTls& PeerPipeTls() { + static PeerTls t; + return t; +} +using hipStream_t = void*; +bool RetirePinThenUnpin(PeerSlot&, hipStream_t, hipStream_t) { return true; } +void RestoreComputeOrThrow(int) {} +bool gate(vt::Queue& compute_q, int slot) { + try { + hipStream_t cst = static_cast(compute_q.handle); + (void)cst; + throw vt::rocm::RestoreFailed{}; + )GATE"; + out << catches; + out.close(); + const std::string cmd = std::string("HIP_VISIBLE_DEVICES= ") + compiler + + " -fsyntax-only -std=c++17 " + path + + " -o /tmp/vllm_prefill_peer_finish_catch_gate.o >/tmp/vllm_prefill_peer_finish_catch_gate.out 2>/tmp/vllm_prefill_peer_finish_catch_gate.err"; + return std::system(cmd.c_str()); +} + +} // namespace + +TEST_CASE("prefill peer source has Launch/Finish and shared cache") { + const std::string hip = ReadHip(); + CHECK(hip.find("LaunchGemma4Fp8ExpertGeGLUPrefillPeer") != std::string::npos); + CHECK(hip.find("FinishGemma4Fp8ExpertGeGLUPrefillPeer") != std::string::npos); + CHECK(hip.find("cache_pin") != std::string::npos); + CHECK(hip.find("PeerSlot") != std::string::npos); + CHECK(hip.find("PrefillDequantCacheT") != std::string::npos); + CHECK(hip.find("ChoosePrefillRetire") != std::string::npos); + CHECK(hip.find("this_gen_ev_e") != std::string::npos); + CHECK(hip.find("RestoreComputeOrThrow") != std::string::npos); + CHECK(hip.find("SameDevLife") != std::string::npos); + const auto prefill = hip.find("bool RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice"); + const auto pinhost = hip.find("void PinGemma4Fp8ExpertHostCache"); + REQUIRE(prefill != std::string::npos); + REQUIRE(pinhost != std::string::npos); + CHECK(hip.find("LaunchGemma4Fp8ExpertGeGLUPrefillPeer", prefill) != std::string::npos); +} + +TEST_CASE("prefill peer Ensure rejects live pin instead of FreeAll") { + vt::rocm::PrefillDequantCacheHost cache; + { + std::lock_guard lk(cache.mu); + REQUIRE(cache.Ensure(0, 4, 8)); + } + void* gu = nullptr; + void* dn = nullptr; + int pin = -1; + const char key = 'k'; + { + std::lock_guard lk(cache.mu); + REQUIRE(cache.GetLocked(&key, &gu, &dn, &pin)); + CHECK(cache.LivePins() == 1); + CHECK_FALSE(cache.Ensure(0, 5, 8)); + CHECK(cache.I == 4); + cache.UnpinLocked(pin); + } + { + std::lock_guard lk(cache.mu); + CHECK(cache.Ensure(0, 5, 8)); + CHECK(cache.I == 5); + } +} + +TEST_CASE("prefill peer pinned slot cannot be rewritten by second worker") { + vt::rocm::PrefillDequantCacheHost cache; + const char ka = 'a'; + const char kb = 'b'; + void* gu = nullptr; + void* dn = nullptr; + int pin_a = -1; + { + std::lock_guard lk(cache.mu); + REQUIRE(cache.Ensure(0, 4, 8)); + REQUIRE(cache.GetLocked(&ka, &gu, &dn, &pin_a)); + int pin_b = -1; + CHECK_FALSE(cache.GetLocked(&kb, &gu, &dn, &pin_b)); + CHECK(cache.slots[pin_a].key == &ka); + cache.UnpinLocked(pin_a); + REQUIRE(cache.GetLocked(&kb, &gu, &dn, &pin_b)); + CHECK(cache.slots[pin_b].key == &kb); + } +} + +TEST_CASE("prefill peer Launch/Finish pairing and fail-after-enqueue retires") { + using vt::rocm::PrefillPeerFailAt; + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + REQUIRE(vt::rocm::HostLaunch(cache, slot, 0, &key, 8, PrefillPeerFailAt::None)); + CHECK(slot.pending_M == 8); + CHECK(slot.cache_pin >= 0); + CHECK(cache.LivePins() == 1); + REQUIRE(vt::rocm::HostFinish(cache, slot, 8, PrefillPeerFailAt::None)); + CHECK(slot.pending_M == 0); + CHECK(cache.LivePins() == 0); + CHECK(slot.compute_restored); + CHECK_FALSE(slot.ev_e_recorded); + CHECK_FALSE(slot.this_gen_ev_e); + + vt::rocm::PrefillPeerLife s2; + REQUIRE_FALSE(vt::rocm::HostLaunch(cache, s2, 0, &key, 8, PrefillPeerFailAt::AfterPinEnqueue)); + CHECK(cache.LivePins() == 0); + CHECK(s2.work_enqueued); + + vt::rocm::PrefillPeerLife s3; + REQUIRE_FALSE(vt::rocm::HostLaunch(cache, s3, 0, &key, 8, PrefillPeerFailAt::RecordEvE)); + CHECK(cache.LivePins() == 0); + + vt::rocm::PrefillPeerLife s3b; + REQUIRE_FALSE(vt::rocm::HostLaunch(cache, s3b, 0, &key, 8, PrefillPeerFailAt::AfterRecord)); + CHECK(cache.LivePins() == 0); + + vt::rocm::PrefillPeerLife s4; + REQUIRE(vt::rocm::HostLaunch(cache, s4, 0, &key, 8, PrefillPeerFailAt::None)); + REQUIRE_FALSE(vt::rocm::HostFinish(cache, s4, 8, PrefillPeerFailAt::AfterWait)); + CHECK(cache.LivePins() == 0); + + vt::rocm::PrefillPeerLife s5; + REQUIRE(vt::rocm::HostLaunch(cache, s5, 0, &key, 8, PrefillPeerFailAt::None)); + REQUIRE_FALSE(vt::rocm::HostFinish(cache, s5, 8, PrefillPeerFailAt::AfterCopy)); + CHECK(cache.LivePins() == 0); +} + +TEST_CASE("prefill peer wrapper uses slot 0 only") { + const std::string hip = ReadHip(); + CHECK(hip.find("LaunchGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0") != std::string::npos); + CHECK(hip.find("FinishGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0") != std::string::npos); + CHECK(hip.find("PeerSlot s[1]") != std::string::npos); + CHECK(hip.find("PeerSlot s[2]") == std::string::npos); + CHECK(hip.find("if (slot != 0 || !x_compute") != std::string::npos); + CHECK(hip.find("if (slot != 0 || !y_compute") != std::string::npos); + CHECK(hip.find("slot < 0 || slot > 1") == std::string::npos); +} + +TEST_CASE("prefill peer failed retirement quarantines pin") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + REQUIRE(vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None)); + REQUIRE_FALSE(vt::rocm::HostFinish(cache, slot, 8, vt::rocm::PrefillPeerFailAt::None, + /*retire_ok=*/false)); + CHECK(slot.quarantined); + CHECK(slot.cache_pin >= 0); + CHECK(cache.LivePins() == 1); +} + +TEST_CASE("prefill peer partial Ensure alloc is freed") { + vt::rocm::PrefillDequantCacheHost cache; + cache.hooks.fail_malloc_at = 2; + { + std::lock_guard lk(cache.mu); + CHECK_FALSE(cache.Ensure(0, 4, 8)); + CHECK(cache.LiveAllocs() == 0); + CHECK(cache.hooks.frees >= 1); + } +} + +TEST_CASE("prefill peer failed ready record keeps fill lease (no cross-stream reuse)") { + vt::rocm::PrefillDequantCacheHost cache; + const char ka = 'a'; + const char kb = 'b'; + void* gu = nullptr; + void* dn = nullptr; + int pin = -1; + { + std::lock_guard lk(cache.mu); + REQUIRE(cache.Ensure(0, 4, 8)); + REQUIRE(cache.GetLocked(&ka, &gu, &dn, &pin)); + cache.UnpinLocked(pin); + cache.hooks.fail_record = true; + int fail_pin = -1; + CHECK_FALSE(cache.GetLocked(&kb, &gu, &dn, &fail_pin)); + CHECK(fail_pin >= 0); + CHECK(cache.slots[0].key == nullptr); + CHECK_FALSE(cache.slots[0].ready); + CHECK(cache.slots[0].filling); + CHECK(cache.slots[0].fill_failed); + CHECK(cache.LivePins() >= 1); + CHECK(cache.hooks.fill_calls >= 1); + int reuse = -1; + CHECK_FALSE(cache.GetLocked(&ka, &gu, &dn, &reuse)); + CHECK(cache.slots[0].key == nullptr); + REQUIRE_FALSE(cache.RetireFillLocked(fail_pin, /*retire_ok=*/false)); + CHECK(cache.slots[0].fill_failed); + CHECK(cache.LivePins() >= 1); + } +} + +TEST_CASE("prefill peer stale prior ev_e is not current rollback target") { + vt::rocm::PrefillPeerLife life; + life.ev_e_recorded = true; // leftover from previous Finish + life.this_gen_ev_e = false; + life.rollback_armed = true; + life.cache_pin = -1; + CHECK(vt::rocm::ChoosePrefillRetire(life) == vt::rocm::PrefillRetireTarget::ComputeStream); + life.this_gen_ev_e = true; + CHECK(vt::rocm::ChoosePrefillRetire(life) == vt::rocm::PrefillRetireTarget::RecordedEvent); + life.this_gen_ev_e = false; + life.rollback_armed = false; + life.work_on_expert = true; + CHECK(vt::rocm::ChoosePrefillRetire(life) == vt::rocm::PrefillRetireTarget::ExpertStream); +} + +TEST_CASE("prefill peer two-invocation: second enqueue fail syncs current stream") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + REQUIRE(vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None)); + REQUIRE(vt::rocm::HostFinish(cache, slot, 8, vt::rocm::PrefillPeerFailAt::None)); + CHECK_FALSE(slot.this_gen_ev_e); + CHECK_FALSE(slot.ev_e_recorded); + REQUIRE_FALSE( + vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::AfterFirstEnqueue)); + CHECK(vt::rocm::ChoosePrefillRetire(slot) != vt::rocm::PrefillRetireTarget::RecordedEvent); + CHECK(slot.work_enqueued); +} + +TEST_CASE("prefill peer restore failure is fatal not false") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + CHECK_THROWS_AS( + vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None, + /*retire_ok=*/true, /*restore_ok=*/false), + vt::rocm::RestoreFailed); + CHECK(slot.pending_M == 0); + CHECK(slot.cache_pin < 0); + CHECK(cache.LivePins() == 0); +} + +TEST_CASE("prefill peer same-dev quarantine blocks reenter and reconfigure") { + vt::rocm::SameDevLife life; + CHECK(life.CanEnter()); + CHECK(life.CanReconfigure()); + life.Quarantine(0, 0); + CHECK_FALSE(life.CanEnter()); + CHECK_FALSE(life.CanReconfigure()); + life.ClearPin(); + CHECK(life.CanEnter()); +} + +TEST_CASE("prefill peer same-dev pin held through GEMM readers") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::SameDevLife life; + vt::rocm::SameDevSession sess; + sess.cache = &cache; + sess.life = &life; + const char key = 'k'; + REQUIRE(sess.Acquire(&key, 0)); + CHECK(cache.LivePins() == 1); + CHECK(life.cache_pin >= 0); + REQUIRE(sess.RunGemmReaders()); + CHECK_FALSE(sess.unpinned_before_gemm); + CHECK(cache.LivePins() == 1); + REQUIRE(sess.Retire(/*sync_ok=*/true)); + CHECK(cache.LivePins() == 0); +} + +TEST_CASE("prefill peer same-dev early unpin before GEMM is RED") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::SameDevLife life; + vt::rocm::SameDevSession sess; + sess.cache = &cache; + sess.life = &life; + const char key = 'k'; + REQUIRE(sess.Acquire(&key, 0)); + { + std::lock_guard lk(cache.mu); + cache.UnpinLocked(sess.pin); // product mutation analogue + } + CHECK_FALSE(sess.RunGemmReaders()); + CHECK(sess.unpinned_before_gemm); +} + +TEST_CASE("prefill peer LifeFromSlot does not remap this_gen to ev_e_recorded") { + struct FakeSlot { + int cache_pin = 3; + int cache_dev = 1; + bool this_gen_ev_e = false; + bool ev_e_recorded = true; + bool rollback_armed = true; + bool quarantined = false; + bool work_on_compute = true; + bool work_on_expert = false; + bool fill_lease = false; + } tls; + const auto life = vt::rocm::LifeFromSlot(tls); + CHECK_FALSE(life.this_gen_ev_e); + CHECK(life.ev_e_recorded); + CHECK(vt::rocm::ChoosePrefillRetire(life) == vt::rocm::PrefillRetireTarget::ComputeStream); +} + +TEST_CASE("prefill peer product RestoreComputeOrThrow no-op mutation is RED") { + const std::string restore = ExtractRestoreComputeOrThrow(ReadHip()); + CHECK(restore.find("if (hipSetDevice(compute_dev) != hipSuccess) throw vt::rocm::RestoreFailed{}") != + std::string::npos); + const auto gxx = CompileAndRunProductRestore(restore, "g++"); + CHECK(gxx.compile == 0); + CHECK(gxx.run == 0); + + const std::string line = + " if (hipSetDevice(compute_dev) != hipSuccess) throw vt::rocm::RestoreFailed{};\n"; + std::string mut = restore; + const auto pos = mut.find(line); + REQUIRE(pos != std::string::npos); + mut.replace(pos, line.size(), " (void)compute_dev;\n"); + const auto mut_gxx = CompileAndRunProductRestore(mut, "g++"); + CHECK(mut_gxx.compile == 0); + // A mutant that fails to COMPILE reads as a caught defect while proving + // nothing, so the compile status is asserted separately from the run status. + CHECK(mut_gxx.run != 0); + + if (const char* hipcc = ResolveHipcc(); hipcc != nullptr) { + const std::string hip_cxx = std::string(hipcc) + " -x c++"; + const auto hip = CompileAndRunProductRestore(restore, hip_cxx); + CHECK(hip.compile == 0); + CHECK(hip.run == 0); + const auto mut_hip = CompileAndRunProductRestore(mut, hip_cxx); + CHECK(mut_hip.compile == 0); + CHECK(mut_hip.run != 0); + } else { + ReportHipccNotRun("RestoreComputeOrThrow no-op mutation, HIP-compiler leg"); + } + + const std::string hdr = ReadHeader(); + CHECK(hdr.find("if (std::uncaught_exceptions() > 0) return;") == std::string::npos); + CHECK(hdr.find("if (!set(dev)) std::terminate();") != std::string::npos); + CHECK(hdr.find("PublishThenRestoreOrThrow") != std::string::npos); +} + +TEST_CASE("prefill peer HostLaunch fill-lease fail retires outside lock") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + cache.hooks.fail_record = true; + REQUIRE_FALSE(vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None)); + CHECK(slot.fill_lease == false); // cleared on successful retire + CHECK(slot.cache_pin < 0); + CHECK(cache.LivePins() == 0); + CHECK_FALSE(cache.slots[0].filling); + CHECK_FALSE(cache.slots[0].fill_failed); + cache.hooks.fail_record = false; + vt::rocm::PrefillPeerLife s2; + REQUIRE(vt::rocm::HostLaunch(cache, s2, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None)); + CHECK(s2.cache_pin >= 0); +} + +TEST_CASE("prefill peer fill-lease failed retire quarantines") { + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + cache.hooks.fail_record = true; + REQUIRE_FALSE(vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None, + /*retire_ok=*/false)); + CHECK(slot.quarantined); + CHECK(slot.cache_pin >= 0); + CHECK(cache.LivePins() >= 1); + const bool lease_held = cache.slots[0].filling || cache.slots[0].fill_failed; + CHECK(lease_held); +} + + +TEST_CASE("prefill peer two-compute-stream output-copy retirement") { + vt::rocm::OutputCopyGate gate; + gate.Enqueue(1); + CHECK_FALSE(gate.CanReuseScratch()); + CHECK_FALSE(gate.Retire(2)); + CHECK_FALSE(gate.CanReuseScratch()); + REQUIRE(gate.Retire(1)); + CHECK(gate.CanReuseScratch()); + + vt::rocm::PrefillDequantCacheHost cache; + vt::rocm::PrefillPeerLife slot; + const char key = 'k'; + REQUIRE(vt::rocm::HostLaunch(cache, slot, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None)); + // Copy on stream 1, retire on stream 2: product mutation analogue is RED. + CHECK_FALSE(vt::rocm::HostFinish(cache, slot, 8, vt::rocm::PrefillPeerFailAt::None, true, true, 1, 2)); + CHECK_FALSE(slot.output_copy.CanReuseScratch()); + CHECK(slot.pending_M == 8); + vt::rocm::PrefillPeerLife reuse; + reuse.output_copy = slot.output_copy; + CHECK_FALSE(vt::rocm::HostLaunch(cache, reuse, 0, &key, 8, vt::rocm::PrefillPeerFailAt::None)); + REQUIRE(vt::rocm::HostFinish(cache, slot, 8, vt::rocm::PrefillPeerFailAt::None, true, true, 1, 1)); + CHECK(slot.output_copy.CanReuseScratch()); + CHECK(slot.pending_M == 0); +} + +TEST_CASE("prefill peer Finish catch HIP/C++ compile gate") { + const std::string finish = ExtractFinish(ReadHip()); + CHECK(CompileFinishCatchTu(finish, "g++") == 0); + std::string mut_cst = finish; + const std::string decl = " hipStream_t cst = static_cast(compute_q.handle);\n"; + auto p1 = mut_cst.find(decl); + REQUIRE(p1 != std::string::npos); + mut_cst.erase(p1, decl.size()); + auto p2 = mut_cst.find(decl); + if (p2 != std::string::npos) mut_cst.erase(p2, decl.size()); + CHECK(CompileFinishCatchTu(mut_cst, "g++") != 0); + + if (const char* hipcc = ResolveHipcc(); hipcc != nullptr) { + const std::string hip_cxx = std::string(hipcc) + " -x c++"; + CHECK(CompileFinishCatchTu(finish, hip_cxx) == 0); + CHECK(CompileFinishCatchTu(mut_cst, hip_cxx) != 0); + } else { + ReportHipccNotRun("Finish catch compile gate, HIP-compiler leg"); + } +} + +TEST_CASE("prefill peer product two-stream output-copy mutation is RED") { + const std::string finish = ExtractFinish(ReadHip()); + CHECK(ProductFinishRetiresCstBeforeReuse(finish)); + const std::string block = + " if (hipStreamSynchronize(cst) != hipSuccess) {\n" + " tls.quarantined = true;\n" + " RestoreComputeOrThrow(compute_dev); return false;\n" + " }\n"; + std::string mut = finish; + const auto pos = mut.find(block); + REQUIRE(pos != std::string::npos); + mut.erase(pos, block.size()); + CHECK_FALSE(ProductFinishRetiresCstBeforeReuse(mut)); +} + +TEST_CASE("prefill peer source: product uses shared retire/restore policy") { + const std::string hip = ReadHip(); + CHECK(hip.find("ChoosePrefillRetire(PeerLifeView") != std::string::npos); + CHECK(hip.find("LifeFromSlot") != std::string::npos); + CHECK(hip.find("RestoreComputeOrThrow") != std::string::npos); + CHECK(hip.find("bool RestoreComputeDevOrFatal") == std::string::npos); + CHECK(hip.find("life.Quarantine") != std::string::npos); + CHECK(hip.find("life.PersistPin") != std::string::npos); + CHECK(hip.find("work_on_compute = true") != std::string::npos); + CHECK(hip.find("FailLaunchRestore(tls, est, cst, compute_dev)") != std::string::npos); + CHECK(hip.find("ReleaseObservedPinLocked") != std::string::npos); + CHECK(hip.find("PublishThenRestoreOrThrow") != std::string::npos); + CHECK(hip.find("tls.fill_lease = true") != std::string::npos); + const auto fin = hip.find("bool FinishGemma4Fp8ExpertGeGLUPrefillPeer"); + REQUIRE(fin != std::string::npos); + const auto fin_def = hip.rfind("bool FinishGemma4Fp8ExpertGeGLUPrefillPeer"); + REQUIRE(fin_def != std::string::npos); + const auto fin_end = hip.find("void PinGemma4Fp8ExpertHostCache", fin_def); + REQUIRE(fin_end != std::string::npos); + const std::string finish = hip.substr(fin_def, fin_end - fin_def); + const auto catch_all = finish.find("catch (...)"); + REQUIRE(catch_all != std::string::npos); + const std::string catch_body = finish.substr(catch_all); + CHECK(catch_body.find("hipStream_t cst = static_cast(compute_q.handle);") != + std::string::npos); + const auto sync_cst = finish.find("hipStreamSynchronize(cst)"); + const auto success_clear = + finish.find("tls.pending_M = 0;\n RestoreComputeOrThrow(compute_dev); return true;"); + REQUIRE(sync_cst != std::string::npos); + REQUIRE(success_clear != std::string::npos); + CHECK(sync_cst < success_clear); + const auto ret = hip.find("bool RetirePinThenUnpin"); + REQUIRE(ret != std::string::npos); + const auto ret_end = hip.find("void FailLaunchRestore", ret); + REQUIRE(ret_end != std::string::npos); + const std::string retire = hip.substr(ret, ret_end - ret); + CHECK(retire.find("(void)hipEventSynchronize") == std::string::npos); + CHECK(retire.find("ChoosePrefillRetire") != std::string::npos); + CHECK(retire.find("PrefillRetireTarget::ComputeStream") != std::string::npos); + CHECK(retire.find("PrefillRetireTarget::ExpertStream") != std::string::npos); + CHECK(retire.find("ReleaseObservedPinLocked") != std::string::npos); + CHECK(retire.find("UnpinLocked(tls.cache_pin)") == std::string::npos); + + // After first successful enqueue is armed, bare restore-return is forbidden. + const auto armed = hip.find("tls.rollback_armed = true"); + REQUIRE(armed != std::string::npos); + const auto rec_end = hip.find("PublishThenRestoreOrThrow", armed); + REQUIRE(rec_end != std::string::npos); + const std::string after = hip.substr(armed, rec_end - armed); + CHECK(after.find("RestoreComputeOrThrow(compute_dev); return false;") == std::string::npos); + CHECK(after.find("FailLaunchRestore(tls, est, cst, compute_dev)") != std::string::npos); + + // Same-dev: no UnpinLocked between GetLocked and MatmulBT. + const auto sd = hip.find("if (expert_dev == compute_dev)"); + REQUIRE(sd != std::string::npos); + const auto launch = hip.find("LaunchGemma4Fp8ExpertGeGLUPrefillPeer(compute_q, /*slot=*/0", sd); + REQUIRE(launch != std::string::npos); + const std::string same = hip.substr(sd, launch - sd); + const auto gl = same.find("GetLocked"); + const auto mm = same.find("MatmulBT"); + REQUIRE(gl != std::string::npos); + REQUIRE(mm != std::string::npos); + const auto unpin = same.find("UnpinLocked", gl); + const bool unpin_after_gemm = unpin == std::string::npos || unpin > mm; + CHECK(unpin_after_gemm); +} + +TEST_CASE("prefill peer this_gen remap mutation is RED") { + const std::string hip = ReadHip(); + CHECK(hip.find("life.this_gen_ev_e = tls.ev_e_recorded") == std::string::npos); + CHECK(hip.find("PeerLifeView(const PeerSlot& tls) { return vt::rocm::LifeFromSlot(tls); }") != + std::string::npos); +}