fix(rocm/gemma4): #838 widen indexed MoE gate to T<=63 + retirement-safe dispatch

.agents/claims/CLAIM-ROCM-GEMMA4-INDEXED-MAX-T.md

@@ -0,0 +1,5 @@
+# CLAIM-ROCM-GEMMA4-INDEXED-MAX-T
+
+| Claim | Row IDs | Agent | Worktree | Branch | Owned scope | State | Last update |
+|---|---|---|---|---|---|---|---|
+| `CLAIM-ROCM-GEMMA4-INDEXED-MAX-T` | `BACKEND-ROCM` (slug `ROCM-GEMMA4-INDEXED-MAX-T`, issue #838) | hermes-vllm (lab), helper | `/home/don/llms/vllm.cpp-indexed-max-t` | `row/ROCM-GEMMA4-INDEXED-MAX-T` | Owns ONLY: widen `gemma4_moe.cpp` T==1 indexed gate to T≤63 via `VT_GEMMA4_DECODE_INDEXED_MAX_T` default 63, using existing per-token indexed helpers, plus tensor oracle vs serial reference. **EXCLUDED:** packed `ExpertGeGLUFp8TopKIndexedBatched`, DEVICE_GROUP, INDEXED_NOSYNC, #837, #839, #697. Independent history from abandoned `row/ROCM-GEMMA4-XDEV-MOE`. | `IMPLEMENTING` | 2026-08-15 — d973/0f32 repair: retire inside acc_idx scope, Retire fail-closed, independent serial ref, injectable arms, owner kind |

.agents/issue-index.md

@@ -511,6 +511,7 @@ rather than merged. `scripts/check-agent-record.py` gates both.
 | [#1493](https://github.com/mudler/vllm.cpp/issues/1493) | `LTX25-DFR-ROUNDS` | **DFR's unclamped `2**round_idx` tile count is mirrored and gated by nothing, because every fixture canvas in this tree has ONE keyframe segment.** `tile_ranges` clamps to `min(num_tiles, n_segments)` (`dfr_layout.py:171`), and the 9-frame fixture pads to a 25-frame canvas with a single segment, so round 1 asks for 2 windows and gets 1 and round 2 asks for 4 and gets 2. A port computing `round_idx + 1`, or `2 * round_idx`, or capping at 2 returns the SAME tile counts on every test here, with every downstream shape, frame count and exit status identical. NOT FIXED IN FLOW and the judgement is recorded rather than reversed: reaching 4 segments needs a materially longer canvas, so round 2 would denoise 4 tiles on a canvas already doubled twice - a new fixture and a substantially longer CPU run in a file that already carries 102 cases, not an assertion added to the existing render. The bound is stated in the test body and in `docs/USAGE.md`, so it was disclosed before it was owned; this row is the ownership. Closing it needs one render whose canvas carries at least 4 segments plus an assertion that `round_tile_counts` reads the unclamped `2**round_idx` for at least one round, which is the only shape that separates `2**round` from every expression agreeing with it at 1 and 2. Listed under `## Owed` in [`ltx25-dfr-rounds.md`](specs/ltx25-dfr-rounds.md) | bug |
 | [#1494](https://github.com/mudler/vllm.cpp/issues/1494) | `LTX25-DEVICE-RESIDENCY` | **A SECOND LTX-2.5 phase-coverage ratio decides by box load, and it is NOT the one [#1439](https://github.com/mudler/vllm.cpp/issues/1439) tracks.** `ltx2 video: the three carrying phases contain their work and the load keeps its order` asserts `CHECK_MESSAGE(covered >= c.min_coverage * leaf_seconds, ...)` at `tests/vllm/multimodal/test_ltx2_video.cpp:3696`; #1439 is `CHECK(leaves >= 0.95 * wall)` in a DIFFERENT case, now at `:3259`. Closing one does not close the other. Measured 2026-08-20, x86_64 `Release` `VLLM_CPP_CUDA=OFF`, three consecutive full-suite runs of ONE binary (`sha256 8fdbc31d...`) with no source change: loadavg 10.45 gives 94.6039% RED, a quieter run gives 96.8506% green, loadavg 16.53 gives 94.6039% RED. The comparison prints as `CHECK( 0.00414483 >= 0.00416218 )` - `denoise` is 0.00438124 s, its eight named sub-scopes cover 0.00414483 s, so the un-named residue is 0.00023641 s and **the margin is a quarter of a millisecond**. Same scheduling polarity #1439 recorded, which is the tell that this is the instrument and not the code: the run that PASSED is the run where `denoise` took 0.00940481 s, more than twice the failing runs', because the residue grows more slowly than the leaf it is divided by. NOT FIXED IN FLOW: bounding the residue in SECONDS beside the ratio, so the assertion says the same thing at fixture and production scale, changes a gate's semantics and needs its own row, spec and red-first evidence per `AGENTS.md` `## Changing the rules or a checker` - the same conclusion #1439 reached, and the two should be repaired together because one seconds bound would serve both. Found by the fresh implementer repairing the review findings of [#1481](https://github.com/mudler/vllm.cpp/pull/1481); pre-existing and not that PR's defect | bug |
 | [#1474](https://github.com/mudler/vllm.cpp/issues/1474) | `VT-CONV1D-F32-ACC` | `vt::Conv1d` and `vt::ConvTranspose1d` accumulate every output cell in **f64** against a reference that accumulates **f32**, and the provenance recorded for that width is wrong in two independent ways. **vLLM owns no such op** at the parity pin `555967922` — no `Conv1dLayer`, no `ConvTranspose*Layer`, and the only `ConvTranspose` string in the tree names `torch.nn` classes inside an NVTX annotation hook (`vllm/utils/nvtx_pytorch_hooks.py:47-49`); vLLM deliberately drops the vocoder it would otherwise own (`qwen3_omni_moe_thinker.py:1975`, `skip_prefixes=["talker.", "code2wav."]`), and where it DOES own a convolution it states this polarity itself (`csrc/cpu/mamba_kernels.hpp` — `// Accumulate in float32 for precision`). So the reference is torch, and torch was MEASURED rather than read: a 27-tap `[+1e8, 0.1 x 25, -1e8]` probe over a uniform input, where an f32 accumulator lands on exactly `0.0` in ANY order and an f64 one on ~2.5, returns **0.0** for `F.conv1d` at f32 AND at bf16, **0.0** for `F.conv_transpose1d` at f32, and 2.4999998510 at f64 (torch 2.11.0+cu130). **The recorded justification fails twice.** `.agents/specs/minimax-music3.md` §13.2 and §18.3 say f64 "is what every committed golden for all four consumers was taken with", but all three generators run torch in f32 — `gen-bigvgan-goldens.py:48` builds f64 then `.float()`, `gen-ltx2-vae-goldens.py:223,234` and `gen-minimax-music3-acoustic-goldens.py:81,134` cast every parameter and input with `astype(np.float32)` — so the goldens are the output of an **f32-accumulating** reference and the f64 arm has been wider than its own oracle since the op landed. And `include/vt/ops.h` clause (1) cites "every committed golden under `tests/parity/goldens/`", a directory whose 101 entries contain **no** vocoder, BigVGAN, LTX-2.5 VAE, FVQ or general-conv1d golden at all; those live in `tests/vllm/models/*_goldens.inc`. An uncheckable citation is how the first claim survived. **The precedent is landed inside one of the same models**: [#1008](https://github.com/mudler/vllm.cpp/issues/1008) narrowed LTX-2.5's conv VIDEO VAE f64->f32 on this exact argument (`d1b0ea3a8`, [PR #1036](https://github.com/mudler/vllm.cpp/pull/1036)), leaving LTX-2.5 shipping f32 video decode and f64 audio decode in ONE model. The measured risk is the summation ORDER, not the width: in #1008 naive-serial f32 pushed one arm to 5.00679e-06 against a 5e-06 tolerance, repaired by per-input-channel BLOCKED partials — which is what torch's blocked-GEMM convolution does and what `vt::Conv3d` already declares as its contract. Spec [`vt-conv1d-f32-accumulator.md`](specs/vt-conv1d-f32-accumulator.md); corrects both false records in the same flow | bug |
+| [#838](https://github.com/mudler/vllm.cpp/issues/838) | `ROCM-GEMMA4-INDEXED-MAX-T` | Gemma4 indexed MoE gate widened to T<=63 with injectable dispatch + retire-in-scope | perf |
 | [#1476](https://github.com/mudler/vllm.cpp/issues/1476) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **TT host-free captured decode went degenerate at the first KV block boundary while host-free eager stayed coherent.** Found by the operator gate at `206afb63`, reproduced at the landed SHA `79ff8f31` in fresh clean builds — not a main regression; the recorded 22/22 argmax predated the final on-device `cur_pos` integration and did not reproduce on the landed tree. TWO root causes, both read at the tt-metal source and fixed together: (1) the RAC `page_table` was allocated `[C,1]`, but `reader_update_cache_interleaved_start_id.cpp:110-111` resolves `page_table_ptr[update_idx / block_size]` — the STICK — so the first write past `cur_pos >= block_size` (32) indexed past the one-column tensor and landed KV in a garbage physical block; the device tensor now carries the user's whole block-table row `[C, block_table_cols]`, grows by reallocation on width change, and RETIRES (keeps alive) the superseded tensor because a freed buffer can hand its address to a new allocation while a recorded trace still addresses it. (2) `WarmDecodePos` skipped its re-seed whenever the process-global `GraphCapturesDone()` was nonzero — a counter `Reset()` never clears — so the cold eager step after a boundary reset ran no plus_one and the RE-captured trace read `cur_pos` one position behind; the regime flag now comes from the driver (`s.graph.captured()`), and every cold/warm/capture step re-seeds. Implementer verification on the P150 (full-answer compares, never a first-line `grep -m1`): the degeneration is GONE — a 160-token captured run is coherent with the 80-token answer as a strict byte-prefix across 5 block boundaries, and reverting either root cause in a scratch build regenerates the word salad (the review's G1/G4 mutations). Captured vs host-free eager is NOT byte-identical — the earlier identical claim was a first-line compare artifact: full answers 284B md5 `3b5a579d82d58396fe4e344826946403` vs 286B md5 `f5ffdf6aa290e11fd187673c2f3c52bb`, first diff at byte 174, both arms coherent. Per-step top-2 adjudication (`VT_TT_DUMP_KV`, the `qwen3-neartie-gap.py` 0.5-nat bar): argmax identical 45/80 steps with top-2 values agreeing to ≤0.5 logits (≤4 bf16 ULP); the first divergence, decode step 46, is a swapped top-2 near-tie — captured `[11:19.75, 311:19.50]` gap 0.25 nats vs eager `[311:19.625, 11:19.50]` gap 0.125 nats, cross-arm deltas 0.125/0.25 logits — inside the near-tie band this repo already tracks for Qwen3-0.6B on TT (#1488 owes the teacher-forced golden re-adjudication); the 34 later argmax differences are prefix divergence, not numeric evidence. The re-seed regime is gate-covered by a new arm: `VT_TT_RECAPTURE_EVERY=8` (9 captures, 71 replays) is byte-identical to the plain captured arm and restoring the old `GraphCapturesDone()>0` early-return reds it; `test_tenstorrent_backend` 23/23 green with the fix, with and without an ambient flag (exit-time segfault pre-existing, [#1486](https://github.com/mudler/vllm.cpp/issues/1486)). Steady state copies the page_table only on content change — zero copies inside a block — so the fix lands the "Phase 2 full" refresh the old comment owed. Spec [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug |
 | [#1486](https://github.com/mudler/vllm.cpp/issues/1486) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`test_tenstorrent_backend` SIGSEGVs after a fully green doctest summary (23/23 cases, 831/831 assertions): exit 139 from static destruction.** The chain is `std::optional<ttnn::Tensor>` destruction into `ttnn::Tensor::deallocate_impl` into `tt::tt_metal::GraphTracker::is_enabled` on a tracker already torn down — the UMD close lines print first. The holders are process-lifetime cache-entry fields at `src/vt/tenstorrent/tenstorrent_ops.cpp:211,396,707,1020,1206`. PROVEN PRE-EXISTING by A/B in one build directory during the #1476 gate: stash the 3-file fix, rebuild, run → 23/23 + exit 139; restore, rebuild, run → identical. Deterministic per run on thalia (aarch64 P150 host). Impact: ctest marks the suite failed while every test passed — the same signature #1476's gate noted once on `test_qwen3_paged_engine`. Repair direction: drain the static caches before device teardown rather than rely on exit order. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug |
 | [#1487](https://github.com/mudler/vllm.cpp/issues/1487) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | **`test_release_metadata` is red on every aarch64 host, so `agent-preflight` cannot go green on the TT dev fleet.** `tests/scripts/test_release_metadata.py:55` stages the HOST arch — `shutil.copy2("/bin/true", ...)` — into the hardcoded `vllm.cpp-0.0.1-linux-x86_64-glibc-cpu.tar.gz` (`:143`); on x86_64 the ELF matches the manifest by accident, on aarch64 the validator correctly refuses: `ELF host architecture does not match manifest: ARM aarch64`. Measured at `206afb63` == `origin/main` with zero local commits and an uncommitted src-only diff; NOT the #1353 disk-full mode (902G free, no scratch-repo failures, the message is the mismatch itself). Repair direction: make the fixture self-consistent (stage an ELF matching the declared arch, or derive the declared arch from the staged ELF) so the validator is exercised for a manifest that lies. Listed under `## Owed` in [`tenstorrent-host-free-forward.md`](specs/tenstorrent-host-free-forward.md) | bug |

.agents/specs/rocm-gemma4-indexed-max-t-donor-gate-911-933.log

@@ -0,0 +1,23 @@
+  // Device-indexed FP8 MoE for small packed T — no router D2H, no host top-k gather.
+  // T=1: original KEEP decode path (hipGraph-stable TLS acc).
+  // T=2..min(MAX, batch_min-1): P1 concurrent-decode gap (research 2026-08-11).
+  // VT_GEMMA4_DECODE_INDEXED_MAX_T: default 63; =1 → T=1 only; clamp [1,63].
+  static const int64_t kDecodeIndexedMaxT = []() -> int64_t {
+    const char* e = std::getenv("VT_GEMMA4_DECODE_INDEXED_MAX_T");
+    if (e == nullptr || !e[0]) return 63;
+    const long v = std::strtol(e, nullptr, 10);
+    if (v < 1) return 1;
+    if (v > 63) return 63;
+    return static_cast<int64_t>(v);
+  }();
+  // P0 overlap enabler (default OFF): keep the enqueue-only indexed path for
+  // prefill T so layer-split does not hipStreamSynchronize after router D2H.
+  // Not a KEEP speed path — per-token launches. Cap 8192 (pipe ubatch).
+  static const int64_t kPrefillIndexedNoSyncMaxT = []() -> int64_t {
+    const char* e = std::getenv("VT_GEMMA4_PREFILL_INDEXED_NOSYNC");
+    if (!e || e[0] != '1') return 0;
+    return 8192;
+  }();
+  const int64_t indexed_max_t =
+      (kPrefillIndexedNoSyncMaxT > kDecodeIndexedMaxT) ? kPrefillIndexedNoSyncMaxT
+                                                      : kDecodeIndexedMaxT;

.agents/specs/rocm-gemma4-indexed-max-t-donor-ok-1088-1110.log

@@ -0,0 +1,23 @@
+    }
+    static std::atomic<int> grp_fail{0};
+    if (grp_fail.fetch_add(1) < 4) {
+      std::fprintf(stderr, "gemma4 moe: device-group FAIL T=%lld — falling back\n",
+                   static_cast<long long>(T));
+      std::fflush(stderr);
+    }
+  }
+  const bool indexed_ok_t =
+      (T >= 1 && T <= indexed_max_t && top_k <= 8 && top_k > 0 && fp8_res &&
+       (kPrefillIndexedNoSyncMaxT > 0 || T < kPrefillBatchMinT));
+  // Indexed only while T is below the group-by-expert crossover (unless nosync).
+  if (indexed_ok_t) {
+    // per-expert scale on device (once per layer/E).
+    struct EscTls {
+      int dev = -1;
+      int64_t E = 0;
+      const void* host_key = nullptr;
+      std::optional<DBuf> sc;
+    };
+    static thread_local EscTls esc_slots[2];
+    EscTls& esc = esc_slots[(compute_dev == 1) ? 1 : 0];
+    float* escale_ptr = nullptr;

.agents/specs/rocm-gemma4-indexed-max-t-donor.md

@@ -0,0 +1,21 @@
+# Donor evidence — #838 indexed-max-T
+
+Pinned **bytes**, not a dirty-tree HEAD. Implementation must copy these slices (or a later
+immutable replacement that research re-reviews), not re-read `/home/don/llms/vllm.cpp`.
+
+Donor includes `kPrefillIndexedNoSyncMaxT` / `DEVICE_GROUP` neighbors. **Those are out of
+scope for this row.** Only `kDecodeIndexedMaxT` (default 63, clamp [1,63]) and the
+`indexed_ok_t` predicate **without** the nosync disjunct are in scope.
+
+| Field | Value |
+|---|---|
+| Donor tree | `/home/don/llms/vllm.cpp` |
+| Donor git HEAD | `2bb4bd8a` (dirty; these slices are **uncommitted** on that tree) |
+| File | `src/vllm/model_executor/models/gemma4_moe.cpp` |
+| Slice A | `rocm-gemma4-indexed-max-t-donor-gate-911-933.log` lines 911–933 SHA256 `d0d28f3d55ff7d526475c9a2a1d028792cc245ba4cc2b421ebf994fed9b96e59` |
+| Slice B | `rocm-gemma4-indexed-max-t-donor-ok-1088-1110.log` lines 1088–1110 SHA256 `5509f3f77dcadd023ce73743e13a2a1a0237d8dd0a8e25a98a765430fc577bd0` |
+| Recipient | `origin/main` `3ce5a1dc` `gemma4_moe.cpp:735` (`T == 1` only) / `:1345` serial |
+| Captured | 2026-08-14 |
+
+`sha256sum` of each slice file must match the table. Do not treat `2bb4bd8a` as a clean
+donor commit.