#2637 [Bug]: gfx1100 2-GPU `--pipeline-windows 2` + MTP: "verify: layer N never rang" (illegal memory access), engine exits on an 18k-token prompt (0.1.44.1, 2x 7900 XTX)

Description

**Setup:** 2× Radeon RX 7900 XTX (gfx1100, 24 GB each) on one box, Strata **v0.1.44.1** (commit `296b59d2`), Ubuntu 26.04.1 LTS, kernel 7.0.0-38, HIP **7.16.26385-0000000** (`amdrocm-core10.1-gfx1100 10.1.0-3`), Intel Core Ultra 7 265K, 93 GiB DDR5. Model Qwen3.8-Flash-Next unsloth **UD-Q4_K_XL**, `--kv int8 --kv-resident 32768 --max-context 200000`.

Exact flags (both runs identical except `--adapt-every`):

```
--serve --pack <pack> --native <gguf> --expert-profile data/expert-profile.bin
--expert-cache auto --prefill auto --short-read 260 --pipeline-windows 2 --spec 4
--mtp <mtp/rt> --max-context 200000 --kv int8 --kv-resident 32768 --vision
--vram-reserve-mib 3072 --pcie-frac 0.00 --spec-min-p 0.70 --pool-workers 13
--adapt-every 1 (run A) / 0 (run B) --adapt-swaps 80 --adapt-decay 0.97
--conversation-cache-mib 8192 --layer-split auto
```

env: `STRATA_PF_GEMM=1`, `STRATA_KV_PREFETCH=1`, `STRATA_HIPBLASLT_TUNING=<gfx1100 tuning file>`. Layer split auto → K=24 (CUDA0 layers 0-23, CUDA1 24-47 + head); the pw2 cost line reports "272 MiB of CUDA0 kept out of the expert cache" (expert cache CUDA1 4933 → 4879 slots).

**Symptom.** With `--pipeline-windows 2` the pipelined decode dies on a **~18k-token prompt** (cold read, `max_tokens` 256, temperature 0):

```
strata serve: pipelined decode: verify: layer 47 never rang (an illegal memory access was encountered)
```

The engine then exits (`exit code 1`) or is reloaded, and the client sees:

| run | request | HTTP | message |
|---|---|---|---|
| A (`--adapt-every 1`) | 16384 cold r1 | 400 | `verify: layer 47 never rang (an illegal memory access was encountered)` |
| A | 16384 cold r2 | 503 | `the engine said nothing for 684 s` |
| B (`--adapt-every 0`) | 16384 cold r1 | 400 | `verify: layer 36 never rang (unspecified launch failure)` |
| B | 16384 cold r2 | 503 | `the engine stopped unexpectedly (exit code 1); the next request...` |

**Turning off the adaptive expert swap changes the symptom, not the outcome.** This was the obvious suspect (a synchronous adaptive swap racing a pipelined speculative window), so the same 18k prompt was run with `--adapt-every 0`: different layer (36 vs 47) and a different HIP error (`unspecified launch failure` 719 vs `illegal memory access` 700), but it still faults. The pipelined decode path itself is the trigger on gfx1100.

**Reproduce rate:** 2/2 runs, 2/2 `--adapt-every` settings, always on the first 16384-token request of the session (real prompt length 18 228 tokens, cold, no cache reuse).

**Small prompts are fine, and the payoff is large** — which is why this is worth fixing rather than just documenting:

| prompt length | result | decode, pw0 → pw2 |
|---|---|---|
| 128 / 512 / 4096 tokens (run A) | clean, 4 passes each | cold 54.5→75.2 / 53.5→80.8 / 52.5→73.5 tok/s; cached 54.5→93.4 / 53.5→82.0 / 52.5→79.6 |
| 128 and 4096 tokens (run B) | clean, 4 passes each | cold tg 87.7 / 81.6; cached tg 85.4 / 76.9 |
| **16384 tokens** (both runs) | **fault** | — |

Prefill is flat (503.8 → 504.9 tok/s at 4096): the gain is entirely decode, up to ×1.7 over pw0 — the serial stage alternation disappearing, as intended.

**Verbatim engine log, run A (`--adapt-every 1`):**

```
strata serve: suffix drafts: 1 windows, 2 of 5 drafts accepted
strata serve: conversation cache: parked 5061 tokens in 140.1 ms; ...
strata serve: pipelined decode: verify: layer 47 never rang (an illegal memory access was encountered)
strata serve: adaptive swap copy failed (layer 7, slot 3984, pinned 0): an illegal memory access was encountered
strata: released the verify window's GPU waits (#267): the GPU finished in 0 ms        (×4)
Hip error: 'an illegal memory access was encountered'(700) at /__w/rockrel/rockrel/rocm-libraries/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp:204
hipModuleUnload failed:  error: an illegal memory access was encountered             (×6)
[strata] ... engine started: ... --adapt-every 1 ...                          <- engine reload
```

**Run B (`--adapt-every 0`):**

```
strata serve: suffix drafts: 5 windows, 19 of 25 drafts accepted
strata serve: conversation cache: parked 5222 tokens in 140.2 ms; ...
strata serve: pipelined decode: verify: layer 36 never rang (unspecified launch failure)
strata: released the verify window's GPU waits (#267): the GPU finished in 0 ms        (×4)
Hip error: 'unspecified launch failure'(719) at /__w/rockrel/rockrel/rocm-libraries/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp:204
hipModuleUnload failed:  error: unspecified launch failure                           (×6)
[strata] ... engine started: ... --adapt-every 0 ...                          <- engine reload
```

The `#267` lines are the watchdog firing *after* the fault — the cause is the verify window never ringing, not the 5 s watchdog.

**Why a gfx1100 box reaches this at all.** `--pipeline-windows 2` is auto-off on gfx1100 (`decode_defaults_arch_ok()` is gfx12-WMMA-only on HIP), but passing the flag explicitly skips that gate: the `pipeline_windows > 0` validation in `generate.cpp` checks serve / layer split / MTP / `--batch` / `--peer-device` / helper caches and never the arch, so the engine starts the pipelined path with no warning that the arch's measured defaults (route tail skip, PCIe share) do not apply. `docs/MULTI_GPU.md` ("`--pipeline-windows 2` asks for it on any split") reads as "supported on any split". Either honour the gate with a refusal/warning on gfx1100, or fix the decode path there.

**Workaround:** `--pipeline-windows 0` restores stability — the identical flag set (pw0) ran the full length sweep including 65536 tokens without any verify fault. The engine is back on pw0 here.

**Related:** #2245 (gfx1030 2-GPU pw2 + MTP → #267 verify *hang*, bisected to `6126f8d`; that one stalls the MEC, this one raises a HIP fault) · #267 · #1341 (closed, >100k prompt deadlocks the pipelined verify-window reader) · #1740 (closed, Windows pw2 stall) · #1447 (closed, docs: pw2 and `--adapt-async 1` combine).

Happy to re-run with a minimal flag set (dropping `--short-read`, the `STRATA_*` env, `--vision`), to bisect the prompt-length boundary between 4.9k and 18.2k, or to try #2245's `STRATA_ATTN_LANECELL=1` on gfx11 if that is useful.