#1705 [Bug]: gfx1100 (RX 7900 XTX): verify timeouts and prefill "no progress" from host-page reclaim; HSA_USERPTR_FOR_PAGED_MEM=0 ends them (#649 mechanism, second GPU)

Descrição

## What happened

Long-context serving on an RX 7900 XTX (gfx1100, Linux) died repeatedly with the two paths this engine already
instruments: the prompt watchdog (`no progress for 60 s … (issue #29)`) and the verify window, whose #649 trace ended in
`RELEASE-NOT-DRAINED`. **Nine engine deaths in one day** (2026-10-08, 10:16 to 23:42), all during long-context work
(positions 31 860 to 200 523 in the trace dumps), all while the box was under memory reclaim. Adding
`"HSA_USERPTR_FOR_PAGED_MEM": "0"` to the server JSON `env` block stopped them: the next long run did **283 requests in
4 h 54 m, up to 245 625 tokens, with no hang, no timeout and no watchdog event**.

This is the mechanism already documented in `docs/AMD_HIP.md` ("Linux: verify timeouts while the kernel reclaims host
memory") and reported for the R9700 in #750 — reported here as a second machine, on a different GPU and ROCm version,
with the extra evidence that the two published workarounds are the same mechanism seen from two ends, and with one
detail that is not in the docs: on this box the trigger was Pop!_OS's default `vm.swappiness=180` zram setup.

Two code-level observations at the end, which are the parts that may want a change in Strata rather than in the config.

### Environment

- Engine **0.1.41** for seven of the nine events (the 0.1.40.3 build hung twice earlier the same day, at 10:16 and
  12:47; the upgrade to 0.1.41 was at 16:40), HIP backend
- **RX 7900 XTX 24 GiB (gfx1100)**, one GPU, 125.7 GiB system RAM
- **Pop!_OS 24.04, kernel 7.1.5-76070105-generic**, ROCm from the pip wheels: `7.10.0a20251120`
  (`_rocm_sdk_core` / `_rocm_sdk_devel` / `_rocm_sdk_libraries_gfx110X_dgpu`), hipBLASLt tuning table
  `tools/hip/gfx1100-hipblaslt-100200.txt`
- Model: Qwen3.8-Flash-Next **UD-IQ4_XS** (unsloth-UD), 48 layers, 12 full-attention (`layer % 4 == 3`), native GGUF in
  place, `--resident-budget-gib 55` (resident RAM mode: 40.65 GiB page-locked experts), `--max-context 262144`,
  `--kv q4_0 --kv-resident 32768`, `--spec 4 --spec-min-p 0.5 --batch-mtp`, `parallel: 2`, `--batch 2`
- `env` at the time of the failures: `STRATA_HIPBLASLT_TUNING`, `GPU_PINNED_MIN_XFER_SIZE=1048576` (already set from
  the #920 note), plus the diagnostics named below

### Symptom — two distinct classes, both fatal

Class A, the prompt path (watchdog stops the engine, exit `-6`):

```
strata serve: no progress for 60 s during a request (reading the prompt (batched): waiting for the GPU (attention,
  router) at layer 9 of the prompt chunk from token 106496) - stopping the engine so the server starts it again (issue #29)
strata serve: stall report (engine 0.1.41): stage "reading the prompt (batched): waiting for the GPU (attention, router)
  at layer 9 of the prompt chunk from token 106496" for 60 s; 0 layers …
  expert pool: epoch 610306, batch epoch 610306: 48 of 48 jobs claimed, 48 done; 15 of 15 workers parked, 15 sleeping
  verify window (last window, not the current stage): 4 tokens at position 192288, host at layer step 48; the GPU rang
  48; flags: served 48, plan (A) 48, copies (B) 48
strata verify trace: GPU breadcrumbs (us after the window's first), the GPU got as far as layer 48:
```

The device had completed all 48 layers (`48 of 48 jobs claimed, 48 done`, `the GPU rang 48`) while the host sat in the
prompt path for 60 s.

Class B, the verify window (the server ends the engine; the console adds that the waits were released and the GPU did not
finish within 5 s):

```
   5000640.1 TIMEOUT              v0x7ffd21865680 window 331 step 4 layer 4 aux 600 | seq 4 flag 4 A 4 B 4
   5000130.7 RELEASE              v0x7ffd21865680 window 331 step -1 layer -1 aux 5000 | seq 4 flag 4 A 4 B 4
         0.0 RELEASE-NOT-DRAINED  v0x7ffd21865680 window 331 step -1 layer -1 aux 5000 | seq 4 flag 4294967295 A … B …
strata verify trace: GPU breadcrumbs (us after the window's first), the GPU got as far as layer 4:
```

Nine events on 2026-10-08 (times are the restart the server did right after each):

| restart after | engine | class | where it stopped |
|---|---|---|---|
| 10:16:24 | 0.1.40.3 | A | stage `(request -1)` |
| 12:47:43 | 0.1.40.3 | A | prefill layer 43, chunk from token 82 501 |
| 17:31:54 | 0.1.41 | A | prefill layer 19, chunk from token 8 192 |
| 18:43:46 | 0.1.41 | A | stage `(request -1)` |
| 19:47:18 | 0.1.41 | A | stage `(decode -1)` |
| 20:56:01 | 0.1.41 | B | window 3761 — **RELEASE-DRAINED** (the GPU finished the rest after release) |
| 22:57:39 | 0.1.41 | B | window 919, layer 10 — RELEASE-NOT-DRAINED |
| 23:21:22 | 0.1.41 | A | prefill layer 9, chunk from token 106 496 |
| 23:48:59 | 0.1.41 | B | window 331, layer 4 — RELEASE-NOT-DRAINED |

The two classes differ in one measurable way, which is what makes them worth separating: after the sentinel release,
class A's GPU work is already done (the host had not seen the ring), class B's never drains (the device was still inside
the step).

### Evidence

1. **Where class B freezes.** Both `RELEASE-NOT-DRAINED` events stopped at the same point in the layer, and the trace
   names it. At 22:57 the same window recorded `out-proj(16)=20628.0 router+ring(…)` for layer 8 and
   `out-proj(16)=22180.5 router+ring(…)` for layer 9, then layer 10 recorded
   `pre(0)=23267.0 … out-proj(16)=23468.0` and **no `router+ring` marker at all** — i.e. it stopped inside the publish of
   that layer's ring to host memory (`verify.cpp:310` names marker 17 `router+ring`). At 23:42 the same: layers 2 and 3
   recorded their markers, layer 4 ended at `out-proj(16)=5636.6` with no `router+ring`. The stall was 23.5 ms and 5.6 ms
   into the window respectively, then nothing for the whole timeout.
2. **The host was not blocked.** A 250 ms sampler over `/proc/<pid>/task/*/stat` (30 MB of dumps) shows the main thread
   `STATE=R WCHAN=0` with ~100 jiffies/turn during every class-A stall: spinning on a ring the device had already
   written, one core saturated, no thread in `D`, `io wait causes` all zero.
3. **Reclaim was happening.** `PSI memory some avg10=5.55%` at the 23:42 hang (0.00% now); the engine reported
   `32 838 major page faults` and `687 MiB in swap`; ~9.4 GiB of the box was in zram swap; `vm.swappiness=180`
   (Pop!_OS default via `/usr/bin/pop-zram-config`).
4. **The driver said so too.** `journalctl -k` carries `amdgpu_amdkfd_restore_userptr_worker 3884: restore worker
   triggers a page restore failed` and `svm_range_restore_pages ... invalid node` in clusters at 10:11, 11:27–11:37,
   12:43 and 17:22 — each before a hang. **Zero since 23:48.** (Caveat: those lines print on count doubling, so silence
   is weak evidence.)
5. **Not context length.** The verify-window dumps of these events carry positions from 31 860 up to 200 523, while 320
   requests in 128–192 k and 8 above 192 k succeeded that day, and the longest clean prompt that evening was 203 908
   tokens. Failure counts track tokens processed, not context size.

## What was ruled out, one variable at a time

Each row is a separate restart of the server with one change, judged on the same workload.

| change | result | what it rules out |
|---|---|---|
| `STRATA_PF_STEP_SYNC=1` (diagnostic) | hangs continued; it named the stalls: 19 585 ms and 8 718 ms prompt steps against a 250 ms log threshold | nothing — but it produced the step names above |
| `STRATA_VERIFY_TRACE=1` (#649 trace) | produced the TIMEOUT / RELEASE breadcrumbs and the layer markers | nothing — it is the instrument |
| ZFS ARC capped to 16 GiB (`zfs_arc_max`) | the 22:57 event was still RELEASE-NOT-DRAINED, `MemAvailable` 87 GiB at the time | ARC growth / ZFS reclaim as the only trigger |
| `STRATA_DOORBELL_STORE=1` | class B returned (23:42) | the read-modify-write ring increment as the whole story — and see observation 2: it cannot reach the class-A path |
| `STRATA_HC_SPLIT=0` (plain hc read) | class B returned (23:42) | the split hc staging path (same finding as #649, which sets this too) |
| `HIP_HOST_COHERENT=1` | class B returned (23:42) | non-coherent host mappings as the mechanism |
| single request, no concurrency | the 18:43 event happened on a single request anyway | `parallel: 2` / `--batch 2` / `--batch-mtp` as a precondition |
| disk I/O left saturated (`PSI io` 63–85% for 18 h from an unrelated `git lfs pull`) | clean for 4 h 54 m under it | block I/O contention as the trigger |
| `HSA_USERPTR_FOR_PAGED_MEM=0` | **clean** | — |

## The fix, and the paired numbers

```json
"env": {
  "STRATA_HIPBLASLT_TUNING": "tools/hip/gfx1100-hipblaslt-100200.txt",
  "GPU_PINNED_MIN_XFER_SIZE": "1048576",
  "HSA_USERPTR_FOR_PAGED_MEM": "0"
}
```

Before (same machine, same model, same args, 17:22–23:42): 7 engine deaths in ~6 h 20 m, longest clean gap 74 min,
prompt reads 614–1 128 tok/s, prompt-step outliers to 19 585 ms.

After (engine 0.1.41, 06:11:21 → 11:05): **4 h 54 m, 283 requests, median prompt 72 242 tokens, largest 245 625,
25 requests above 150 k, zero hangs, zero timeouts, zero prompt-step outliers, zero engine deaths.** Four cold full
prompts in that window: 203 272 tokens in 209 982 ms (968 tok/s), 204 295 in 211 431 ms (966), 205 311 in 211 132 ms
(972), 212 721 in 229 558 ms (927). Decode 34.8–40.4 tok/s with `--spec 4` drafts accepted (e.g. 266 of 320), which
matches the pre-change decode rate. `PSI memory` 0.00%, kernel userptr-restore warnings 0.

Two honest caveats about attribution:

1. **Two changes landed 12 seconds apart.** `swapoff -a` at 23:48:28 and the `env` edit, with the engine starting at
   23:48:58. Both act on the same mechanism — the kernel moving host pages the GPU can see — so the finding is
   "one mechanism, two mitigations", not "switch X". Swap was re-enabled at 11:49 with `vm.swappiness=20` and the
   switch left in; that run was clean for its first 29 minutes, which is the useful configuration and still short.
2. **Speed parity was not measured as a pair.** The switch has not been removed since, so I cannot quote a
   with/against number like the #750 report does; the 927–972 tok/s above is a before/after comparison across
   different prompt mixes, not an interleaved A/B.

A structural change I can see from userland, which supports the mechanism: with the switch on, the engine's `VmRSS`
fell from 78–94 GiB to 6.3 GiB while the log still reports `resident RAM mode: 40.65 GiB of experts in RAM
(page-locked)` and the decode hit rate stays 84–89.5%, and `Unevictable`/`Mlocked` stay at 8 MiB. The 40 GiB appears to
be held outside the process's page tables now (about 57 GiB of `used` memory is not attributed to any process or to
`/proc/meminfo` categories), i.e. it is no longer paged host memory the kernel may move. That is an inference from
accounting, not a driver counter.

## Two things in Strata that look like bugs

1. **`death_note()` only scans the last 4096 bytes of the log** (`serve/server.py:781-796`). With
   `STRATA_VERIFY_TRACE=1` the trace dumps are ~30 KB, so the `issue #29` line is pushed out of the tail and the server
   reports the fallback "running out of RAM" note for what was a caught hang. It sent me chasing memory for an hour.
   Either scan further back or keep the last watchdog line as state.
2. **`STRATA_DOORBELL_STORE` only covers the verify path.** `g_doorbell_store` is read in `src/core/verify.cpp:130` and
   applied at `:1339`, but the prompt path publishes through `doorbell_ring()` (`src/core/layer.cpp:391` →
   `src/kernels/cuda/elementwise.cu:428`, kernel at `:209`), which is an unconditional increment. So the A/B switch
   cannot test the class-A path at all, and the banner ("doorbell stored") reads as if it had. The comment at
   `elementwise.cu:189` about the cloned-literal host expression suggests the same fix applies there.

## What I did not test

- No second machine, no other ROCm version, no other card; one box, one GPU, one model.
- No controlled speed A/B for the switch (see caveat 2), and no output-equality check like the #750 report's five paired
  runs — I did not veri