fix(rocm/gemma4): #839 retirement-safe prefill peer GeGLU Launch/Finish

[bakon11]

fix(rocm/gemma4): #839 retirement-safe prefill peer GeGLU Launch/Finish

Fixes #839.

The defect

The Gemma-4 FP8 cross-device prefill helper was one monolithic function. Its
donor unpinned the dequant cache immediately after enqueueing the GEMMs and
before hipEventRecord(ev_e), and the cache is process-wide, so a second worker
could evict a zero-pin slot while the first expert stream was still reading it.
Under batched-MoE prefill that is a use-after-free, and the observed symptom is a
wedge after matched BEGIN/END with the process idle in kfd_wait_on_events.

What changed

RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice stays the public symbol and
becomes a wrapper over LaunchGemma4Fp8ExpertGeGLUPrefillPeer and
FinishGemma4Fp8ExpertGeGLUPrefillPeer. The cache pin now lives on the
PeerSlot and is released only after the host has observed retirement, never
after a mere enqueue:

• Launch pins under the same mutex scope as GetLocked and stores
  {cache_pin, cache_dev} on the slot before any GEMM is enqueued. It never
  unpins.
• Finish host-waits ev_e, then hipStreamSynchronize on the compute stream
  that received the output copy, and only then unpins under the lock.
• A failed fill or ready-event keeps a fill lease; the lease is retired with
  RetireFillLocked after the producer-stream sync and never under cache.mu.
• A restore failure after publish retires before it rethrows; a failed retire
  quarantines the pin rather than dropping it.
• The same-dev arm follows the same rule. "Enqueue then unpin" is forbidden on
  both arms, because a hipStreamWaitEvent on the compute stream is not a
  host-side retirement proof.

Peer-pipe overlap stays default OFF, so Launch and Finish run back-to-back on
slot 0 and the event order is the one shipping today plus slot-scoped scratch.

What the tests prove

tests/vt/test_ops_gemma4_prefill_peer.cpp is 23 cases over the host lifetime
model and the product source. The lifetime cases drive HostLaunch/HostFinish;
the product is bound to them by source-slice gates that extract the real
Finish and RestoreComputeOrThrow bodies out of
src/vt/rocm/rocm_gemma4_experts.hip, compile them, run them, and require each
mutation to red. Replacing the hipSetDevice(compute_dev) throw in the product
RestoreComputeOrThrow with (void)compute_dev; fails three assertions.

CI repairs in this revision

Three checks were red for reasons that had nothing to do with the change under
review, and all three are green on main.

build-test-cpu and both sanitize-cpu legs failed because the test named one
absolute HIP compiler from the contributor's box,
/opt/rocm-7.2.4/core-7.14/bin/hipcc, in four places and CHECKed that it
returned 0 — so any machine without that exact path failed rather than
reporting that the gate had not run. The toolchain is now resolved
(VLLM_CPP_HIPCC, HIPCC, ROCM_PATH/bin/hipcc, hipcc on PATH,
/opt/rocm/bin/hipcc, each probed with --version) the way
tests/vt/test_ops_getblas_product.cpp:15-21 resolves its own precondition, and
its absence prints a loud HIP COMPILE GATE NOT RUN banner on stderr instead of
asserting. Measured on this tree: 23/23 cases and 199 assertions with no
toolchain plus two banners, 23/23 and 205 assertions with VLLM_CPP_HIPCC set,
the six recovered assertions being exactly the HIP legs. Both arms still red on
the RestoreComputeOrThrow mutation, so the repair stopped the gate failing on
absent hardware without widening it.

Absent hipcc does not exit(77) the way test_ops_getblas_product.cpp does.
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, and exiting
would take those with it. The g++ leg of both compile gates still runs, so
neither case can report a zero-assertion pass.

pr-size refused the whole change with could not classify because nothing in
scripts/check-pr-size.py matches .agents/evidence/. The five donor slices
move beside the spec as SPEC_EVIDENCE .log files (.txt is unclassified
too) and the manifest becomes
.agents/specs/rocm-gemma4-prefill-peer-helper-donor.md. Renaming does not touch
bytes and all five SHA256 values in the table verify unchanged.

docs/FEATURES.md returns to main's text. The branch had rewritten that row
and lost the spec link, the VT_GEMMA4_*/VT_ATTN_* pointer and the
test_gemma4_rocm_fp8_seams seam name in order to fit; main's cell is already
219 of the 220 characters check-public-doc-tables allows. Nothing here owes
that page: check-doc-checkpoint has keyed feature_surface off a changed
REGISTER_VLLM_MODEL set since #595, and this change touches no file under
src/vllm/model_executor/models/.

The branch was rebuilt by rebase onto affc2a7fd, so it carries no untrailered
merge commit and all ten original commits keep their authorship. The
docs/USAGE.md conflict against #837's landed GetBlas text was resolved as a
union: both paragraphs survive.

What this does NOT close

Named here and under ## Owed in
.agents/specs/rocm-gemma4-prefill-peer-helper.md. Row
ROCM-GEMMA4-PREFILL-PEER-HELPER owns them and #839 stays open until they are
discharged.

• No measurement for the new blocking retirement. Finish now host-waits
  hipEventSynchronize(ev_e) and hipStreamSynchronize(cst) on every call, and
  the same-dev arm gains one too, where the path was fully asynchronous before.
  It runs per expert per layer during prefill, so it serialises a pipeline that
  used to overlap. The wait is required for correctness and correctness comes
  first, so it stays; what is missing is the number. Prefill throughput before
  and after on 2x R9700 needs the hardware, and no CI runner here has it.
• PeerSlot s[2] with only slot 0 reachable from the production wrapper.
• The host simulator is a second implementation. HostLaunch/HostFinish in
  include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h are hand-written
  analogues of the product Launch/Finish, ~250 of that header's lines, compiled
  into every HIP build. Moving them under tests/ is the fix; it was not done
  here because only a ROCm box compiles the one translation unit that consumes
  the header, so the move cannot be verified where this repair was made.

Known-unrelated CI

windows-msvc-cpu and windows-msvc-vulkan are red on every open PR from a
break predating this branch (#503, #584).

FOLLOWING_AGENTS_PROTOCOL

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

[localai-bot]

Sending this one back — the lifetime state machine is carefully thought through,
but three things need to change before it can land, and two of them are not
things we can repair for you.

1. The test hardcodes your toolchain path, so it fails everywhere else.

tests/vt/test_ops_gemma4_prefill_peer.cpp:438, :452, :522 and :531 each
pass "/opt/rocm-7.2.4/core-7.14/bin/hipcc -x c++" and then CHECK that the
compile returns 0. On any machine without that exact path it does not skip, it
fails. tests/CMakeLists.txt:1440 registers the suite unconditionally, so
this is the whole of the build-test-cpu, sanitize-cpu (address,undefined)
and sanitize-cpu (thread) red — one suite, 363 - test_ops_gemma4_prefill_peer,
23 cases / 21 passed / 2 failed. All three of those jobs are green on main, so
unlike the windows-msvc-* reds this one is genuinely attributable.

The minimum fix is to resolve hipcc from HIP_PATH/ROCM_PATH/PATH and
skip when absent — tests/vt/test_ops_getblas_product.cpp:20 already does
exactly that with std::exit(77), and your own #1045 uses that pattern.

2. Underneath that is a structural problem the path fix does not solve.

The suite shells out to a compiler at ctest time (std::system at :96, :98,
:158) and greps production source through VLLM_CPP_SOURCE_DIR (:17-27, used
by 8 of 23 cases). Nothing else in this tree tests that way.

More importantly,
include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h:346-620 — roughly 250 of
the header's 624 lines — is a host simulator compiled into every HIP build, and
HostLaunch/HostFinish are a hand-written second implementation of the
production Launch/Finish logic. The tests exercise that copy; the only thing
tying it to the product is the source greps above. So a defect in the real
Launch would not move these tests. AGENTS.md is explicit on both halves
("never write a parallel path by hand", and a unit test that constructs the type
by hand proves the class works, never that anything reaches it). The shared state
machine — PrefillPeerLife, ChoosePrefillRetire, PrefillDequantCacheT,
ComputeDevGuard, RestoreFailed, SameDevLife, OutputCopyGate,
SlotReusable — should stay in the header; the Host* half belongs under
tests/.

3. The perf question the change raises, on a perf issue, is unanswered.

FinishGemma4Fp8ExpertGeGLUPrefillPeer now does a blocking
hipEventSynchronize(tls.ev_e) and a blocking hipStreamSynchronize(cst) on
every call, and the same-dev arm gains a hipStreamSynchronize(cst) too. Before
this change the function was fully asynchronous — record event,
hipStreamWaitEvent, async peer copy, return. It runs per expert per layer
during prefill, so this converts an overlapped pipeline into a host-serialized
one. Nothing in docs/BENCHMARKS.md, the spec, or the PR body records a prefill
throughput or TTFT number either side of it. We have no gfx1201 here, so this
one needs you.

A cross-PR landmine, which decides the order of your stack.

tests/vllm/models/test_gemma4_indexed_max_t.cpp:299-306 (from #1046) slices
rocm_gemma4_experts.hip between "bool RetireGemma4Fp8TopKIndexedPeer" and the
next "RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice", and asserts return true;
is absent from that gap. This PR inserts ~160 lines into exactly that gap. I
extracted the merged tree and confirmed the segment then contains return true;.
That suite is registered unconditionally and reads the .hip from disk, so it
runs on build-test-cpu — meaning #1047 landing after #1046 turns main
red. Best fixed in #1046 by anchoring that assertion to the function body rather
than to the gap between two symbols.

Smaller things, none blocking on their own

• PeerSlot s[2] with only slot 0 ever passed from production. docs/USAGE.md
  says "Peer-pipe overlap stays off (slot 0 only)", but the Nothing lands dead
  rule wants the commit body and PR body to name the owning row and issue, and
  the spec to list it under ## Owed. The spec has no ## Owed section.
• docs/FEATURES.md:352 — the rewrite drops the spec link, the VT_GEMMA4_* /
  VT_ATTN_* env pointer and the test_gemma4_rocm_fp8_seams seam name.
  Appending rather than replacing keeps those.
• .agents/evidence/rocm-gemma4-prefill-peer/MANIFEST.md is an unclassified path,
  which is the pr-size red: ERROR: PR size check could not classify the change. Renaming under .agents/specs/ classifies (SPEC_EVIDENCE matches
  .agents/specs/*.{patch,json,log}). Please do not patch check-pr-size.py —
  that is a governance checker and owes its own spec and mutation evidence.
• The PR body needs the FOLLOWING_AGENTS_PROTOCOL paragraph and the
  Following-Agents-Protocol / AI-Assisted / Assisted-by trailers. The repo
  squashes with squash_merge_commit_message = PR_BODY, so the body is the
  landed commit message. Your branch commits carry them; the body does not.
• windows-msvc-cpu / windows-msvc-vulkan are not yours — they were red on
  every open PR from a break that predates this branch. That queue has since been
  worked down (windows-msvc-cpu/vulkan red on every PR based on c7cb59fbb: C4244 double->float in ltx2_video.cpp's StreamState<->Ltx2LatentState positions copy (#964) #968 landed as fix(ltx2): cast positions double/float to avoid MSVC C4244 (#968) #983, main does not compile under MSVC: #1054 removed a lambda capture that was not redundant #1068 as fix(qwen3.5): the capture #1054 dropped was not redundant, and MSVC says so #1069) and only test_openai_api_server crashes on Windows with STATUS_STACK_BUFFER_OVERRUN (0xC0000409), unmasked by the #512 fix #584's runtime crash
  remains.

One last note that applies to all three of your open PRs: the Verification
section is identical in #1045, #1046 and #1047 and describes a single combined
#837+#838+#839 smoke on d1b0ea3a. Landing any one alone therefore lands
something never exercised in isolation. #1045 is in good shape and close to
landing; splitting your evidence per PR would let it go first.

[localai-bot]

@bakon11 — I pushed three repairs to fix/gemma4-839-prefill-peer (3df512a11 → cf31e5e98, rebased onto affc2a7fd, all ten of your commits keep their authorship, no merge commit). Everything I could fix without an AMD box is fixed. Two things are left, and both need your hardware.

The three CI reds were not about your change. All three are green on main.

build-test-cpu and both sanitize-cpu legs: tests/vt/test_ops_gemma4_prefill_peer.cpp named /opt/rocm-7.2.4/core-7.14/bin/hipcc in four places and CHECKed that the compile returned 0, so any machine without that exact path failed rather than reporting that the gate had not run. The toolchain is now resolved — VLLM_CPP_HIPCC, HIPCC, ROCM_PATH/bin/hipcc, hipcc on PATH, /opt/rocm/bin/hipcc, each probed with --version — the way tests/vt/test_ops_getblas_product.cpp:15-21 resolves its own precondition, and its absence prints a loud HIP COMPILE GATE NOT RUN banner on stderr instead of asserting. Measured here: 23/23 cases and 199 assertions with no toolchain plus two banners; 23/23 and 205 with VLLM_CPP_HIPCC set, the six recovered assertions being exactly the HIP legs. Both arms still red on the RestoreComputeOrThrow mutation (3 failed assertions), so the gate did not get weaker — it stopped failing on absent hardware.

I did not exit(77) the way getblas does. That file is one HIP gate and nothing else; yours carries 21 host-lifetime cases that are the only gate CI has on this change, and exiting would take them with it. If you would rather have the HIP leg as its own skippable CTest entry, that is a clean follow-up — it needs the helpers moved into a tests/support/ header, which I did not want to do blind.

One thing worth knowing for next time: at lines 454 and 531 the old test asserted mut_hipcc.run != 0 after the compile had already failed, so run was 127 and those two assertions passed vacuously. A mutant that fails to build reads as a caught defect. The repaired version asserts the compile status separately.

pr-size refused the whole PR with could not classify because nothing in scripts/check-pr-size.py matches .agents/evidence/. The five donor slices now sit beside the spec as SPEC_EVIDENCE .log files (.txt is unclassified too, so .log rather than a checker change), and MANIFEST.md became .agents/specs/rocm-gemma4-prefill-peer-helper-donor.md. Renaming does not touch bytes — I re-ran sha256sum and all five values in your table verify unchanged.

docs/FEATURES.md is back to main's text. Your rewrite dropped the spec link's companions — the VT_GEMMA4_*/VT_ATTN_* pointer and the test_gemma4_rocm_fp8_seams seam name — to make room, and main's cell is already 219 of the 220 characters check-public-doc-tables allows, so it cannot absorb a sentence. Nothing here owes that page anyway: check-doc-checkpoint has keyed feature_surface off a changed REGISTER_VLLM_MODEL set since #595, and this PR touches no file under src/vllm/model_executor/models/.

The PR body now carries the trailer block, which is what commit-protocol-tag was failing on — your commits always had it.

---

What still needs you.

1. The blocking retirement has no measurement. FinishGemma4Fp8ExpertGeGLUPrefillPeer now host-waits hipEventSynchronize(tls.ev_e) and hipStreamSynchronize(cst) on every call, and the same-dev arm gains a sync too, where the path was fully asynchronous before. It runs per expert per layer during prefill, so it serialises a pipeline that used to overlap — on a perf-labelled issue with no number in docs/BENCHMARKS.md, the spec or the body.

To be clear about where I stand: I am not asking you to remove the wait. Design rule 2 in your own spec is right, correctness comes first, and "enqueue then unpin" genuinely is a use-after-free against a process-wide cache. What is owed is the cost, not a rollback. Prefill throughput before and after on the 2x R9700 at the T values you smoked this on, and if it is expensive, that becomes the next row rather than a reason to ship the race.

2. PeerSlot s[2] with only slot 0 reachable from RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice. Either a production path that reaches slot 1, or narrow the array. Recorded under ## Owed for now so it is visible debt rather than a surprise later.

I also recorded a third item under ## Owed that is on us, not you: HostLaunch/HostFinish in include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h are a hand-written second implementation of the product Launch/Finish, about 250 of that header's 624 lines, compiled into every HIP build — so the lifetime cases exercise the analogue, not the product. What binds them today is your source-slice mutation gates, which is why this is debt and not a blocker. Moving the Host* half under tests/ is the fix; I left it alone because only your box compiles the one TU that consumes that header, so I could not verify the move, and it would churn your other branches' includes.

Ordering with #1046, measured on a real merged tree: your FunctionBody() brace-matcher does its job. On the merge of the two rebased branches, src/vt/rocm/rocm_gemma4_experts.hip auto-merges clean, the FunctionBody slice of RetireGemma4Fp8TopKIndexedPeer is 13 lines on both trees and contains no return true;, and the whole #838 suite is 10/10, 1101 assertions against the merged sources. The old symbol-to-symbol slice would have grown 17 → 177 lines and picked up a return true;, so the de-brittling was load-bearing, not luck. Nice catch.

[bakon11]

Accepted — thanks for the CI toolchain-resolution repairs (and the vacuous-mutation catch at 454/531). Queued on our box, landing on this branch on top of cf31e5e98: (1) the blocking-retirement cost measurement (before/after prefill throughput at the smoked T values — agreed the wait stays; the cost gets recorded, and if it's big it becomes the next row), (2) the PeerSlot s[2] slot-1 answer (wire or narrow). The Host* analogue-under-tests move we'll take as follow-up debt since only our box compiles that TU.