Issues / #1996

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

open · @mantovaniluca91 · 0 comments · View on GitHub

BenchmarksSetup & installServer & APIMulti-GPUAMD / HIPNVIDIA / CUDAModels & quants

Description

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

Related on strata.com

Editorial links to help you install, pick models, or read release notes — not part of the upstream thread.