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

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

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

.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];
+}

.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<size_t>(2 * i_dim) * static_cast<size_t>(h_dim) * 2;
+    const size_t dn_b = static_cast<size_t>(h_dim) * static_cast<size_t>(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<const void*> uniq;
+      static int uk_logs = 0;
+      std::lock_guard<std::mutex> 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<unsigned long long>(misses),
+                     static_cast<unsigned long long>(hits), uniq.size(), key);
+        ++uk_logs;
+      }
+    }
+    return true;