#1996 0.1.43 fixes the #1712 reproduction, but Memory access fault after 31 min under multi-slot load with image requests (2x RX 7900 XTX)

说明

**Strata version, GPU, OS:** 0.1.43 (built from the tag, `STRATA_ISA_FLOOR=avx`) · 2x RX 7900 XTX 24GB (gfx1100), ROCm 7.2.4 in Docker · Ubuntu 22.04.5, kernel 6.8, Xeon E5-2630 v2 (AVX only)

Follow-up to #1712, which 0.1.43 closed. Thanks for taking the fix so quickly. We ran 0.1.43 on the same box the same evening.
Short version: it removes the deterministic image→text fault from #1712, but under our 4-hour multi-slot load it still faulted, after
31 minutes. Patch C on 0.1.41 (attached to my last comment in #1712) passed that same load for 4 hours.

Setup for both runs below: 2x RX 7900 XTX (gfx1100), ROCm 7.2.4 in the container, Xeon E5-2630 v2 (AVX only), 0.1.43 built from the
tag with `STRATA_ISA_FLOOR=avx`, no patches of ours (checked: no `[floria patch]` line, `session is up (engine 0.1.43)`).
Config: `parallel 4`, `--max-context 262144 --kv int8 --kv-resident 32768 --spec 4 --mtp --no-prefill-borrow --batch-groups auto
--trim-stage-weights --layer-split 27 --conversation-cache-mib 6144 --conversation-cache-slots 4`, vision encoder on a separate card
(`cuda_device`).

## 1. The image→text reproduction: fixed

The same two requests as in #1712: a vision request (1803 prompt tokens) and right after it the first text-only
request (2142 tokens, two 2048 chunks). Fresh engine for every run, debug agent attached, 5 alternating rounds, with the order of the
arms reversed in even rounds:

| Engine | Faults |
|---|---|
| 0.1.41 as shipped | **5 of 5** (3 named by the debug agent: `block_topk_kernel`, `MEMORY_VIOLATION`, as before; in the other 2 the debug agent hit its own `wave ... not found` error at the same point) |
| 0.1.43 | **0 of 5** |
| 0.1.43 with `STRATA_NO_IQ128=1` | **0 of 5** |

The third arm turns off the new AVX1 expert kernels. They change the CPU-side timing, and this fault was timing-sensitive. Holding
there too says the upload fix is what removes it.

Also, since the 0.1.42 notes say the AVX1 kernels were only tested under emulation: they run on a real Ivy Bridge here
(`the expert kernels run on AVX1 128-bit (iq_avx1, the older-CPU build)`). They were on in 7 engine starts (5 in the reproduction, 2 in the stress
below), with no SIGILL.

## 2. Four-hour stress: 0.1.43 faults after 31 minutes

This is the same stress that patch C passed on 0.1.41 (240 min, 0 faults), with the same config (including `STRATA_HIP_WMMA=1
STRATA_SELECT_WMMA=1`, as in the C run), and no debug agent. The debug agent restarts the engine under load (`rocm-dbgapi: fatal error: wave_... not
found`), so we leave it out of long runs.

**Load:** 6 clients on the 4 slots, so conversations keep moving through the conversation cache:
- 2 agent sessions on real read-only requests;
- 2 delivery-note loops (vision request, then a structured text request), 4 minutes apart;
- 2 clients that pick at random between a new 32K/64K/128K prompt, a synthetic image followed by a long synthetic text, and a short
  chat.

**Result:** one `Memory access fault by GPU node-4 ... Reason: Page not present or supervisor privilege` at minute 31, and the engine
exited with code -6. No `upload failed` line appeared before it. These are the last engine lines before the fault (hit-rate and
KV lines trimmed). The 1803-token prompts have the size of the delivery-note vision request and are read in the batch while the other
slots decode; then come two short text prompts:

```
strata batch: slot 1 takes 1803 tokens (copied in 84.4 ms)
strata serve: prompt 1803 tokens = 0 reused + 1803 read in 2391 ms (754.1 tok/s), 1 generated in 119 ms (8.4 tok/s), drafts accepted 0 of 0, 1 checkpoints
strata serve: conversation cache: parked 1803 tokens in 85.3 ms; parked=4 bytes=2751082592 evictions=85 snapshot_bytes=264433456 reused_kv_bytes=0
strata batch: slot 3 takes 1803 tokens (copied in 70.3 ms)
strata serve: prompt 1803 tokens = 0 reused + 1803 read in 2086 ms (864.4 tok/s), 1 generated in 30 ms (32.9 tok/s), drafts accepted 0 of 0, 1 checkpoints
strata serve: conversation cache: parked 1803 tokens in 88.6 ms; parked=4 bytes=2187382416 evictions=86 snapshot_bytes=264433456 reused_kv_bytes=0
strata batch: slot 0 takes 70 tokens (copied in 97.2 ms)
strata serve: prompt 70 tokens = 0 reused + 70 read in 1055 ms (66.3 tok/s), 1 generated in 34 ms (29.1 tok/s), drafts accepted 0 of 0, 1 checkpoints
strata batch: slot 2 takes 70 tokens (copied in 104.6 ms)
strata serve: prompt 70 tokens = 0 reused + 70 read in 717 ms (97.7 tok/s), 1 generated in 28 ms (36.2 tok/s), drafts accepted 0 of 0, 1 checkpoints
strata batch: slot 1 takes 1571 tokens (copied in 118.7 ms)
strata serve: prompt 1571 tokens = 0 reused + 1571 read in 2081 ms (754.9 tok/s), 1 generated in 29 ms (34.5 tok/s), drafts accepted 0 of 0, 1 checkpoints
strata serve: KV streaming: 85.69% of 16482 block reads hit VRAM, 9.5 MiB read from RAM
strata serve: conversation cache: parked 1571 tokens in 84.6 ms; parked=4 bytes=1001258752 evictions=89 snapshot_bytes=260888432 reused_kv_bytes=0
Memory access fault by GPU node-4 (Agent handle: 0x8efbf50) on address 0xc9f2d000. Reason: Page not present or supervisor privilege.
```

Without the debug agent we have no kernel name for this one. Until then, the quality and speed matched patch C: one full
delivery-note pass scored 32/37 at 114 s per note under load, against C's 32-33/37 at about 110 s.

## 3. What C covers that 0.1.43 does not

In 0.1.43 the step records and the batched token ids go up from pinned memory. Patch C also moves every other async H2D of the prompt
path whose host side is pageable through pinned bounce slots, one event per slot. In 0.1.43 those copies are:
- the image embedding rows, `cudaMemcpyAsync(m.emb + t * N, row, ...)` from `embd_rows`, at prefill.cpp:2185;
- the draft bounds and tokens in the draft K/V step, at prefill.cpp:1447 and 1451 (`bh[4]` on the stack, `tk` a vector);
- the PLE rows in the `ple_pageable` fallback, at prefill.cpp:2212 (only taken without pinnable RAM);
- the expert routing copies in the non-mapped branches. These are inactive with the default config, where `grp_host` is mapped.

The fault came right after vision requests read in the batch, so the image-row copy is our first suspect. We have not isolated it,
though: C changed all of them at once, and that is the version that held 4 hours.

Would you extend `PinnedUpload` to these copies, or take the bounce-slot approach of C (the patch is attached to my last comment in #1712)? We are keeping
patch C in production for now, and we can run any build of yours through the same reproduction and stress.