Repository navigation
[gfx1200, 2x RX 9060 XT layer split] prompt-read stalls (watchdog #29 aborts) + run-to-run prefill degradation: fresh prompts slow 3-4x mid-run, file tier re-reads 24-35 GB per request #1407
Description
[gfx1200, 2x RX 9060 XT layer split] prompt-read stalls (watchdog #29 aborts) + run-to-run prefill degradation: fresh prompts slow 3-4x mid-run, file tier re-reads 24-35 GB per request
Environment
| Engine | v0.1.40.2, self-built, backend: hip, archs: [gfx1200] |
| ROCm | 10.1.0-3 (gfx1200), HIP 7.16.26385, clang 24.0.0git |
| GPUs | 2x RX 9060 XT 16 GB, --layer-split auto (0-23 / 24-47) |
| Host | Ryzen 5 4600G (6C/12T), 32 GB DDR4, no swap, kernel 7.0.0-38 |
| Model | Qwen3.8-Flash-Next IQ2_XS, ~39 GB experts, --mmap-experts |
| hipBLASLt | self-tuned gfx1200-hipblaslt-100401.txt, log confirms tuning enabled (26 rows, version 100401) — table pairing is not the issue |
Flags: --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --mtp
- Prompt-read stalls (Critical bug + Performance report RTX 4090 #29 watchdog abort)
Fresh 4K-38K prompts: the progress line freezes mid-chunk, the watchdog kills the engine:
[strata] reading the prompt: 16,384 of 38,830 tokens, 495 s so far <- unchanged for 100+ s
strata serve: no progress for 60 s during a request (reading the prompt (batched):
layer 34 of the prompt chunk from token 0) ... (issue #29) -> exit -6
While frozen: both GPUs at 100% busy, 3 threads in D state at lock_mm_and_find_vma, 2 at kfd_wait_on_events, MemAvailable ~17 GB (no OOM). One run also hit verify: timed out at layer 8 (#267); an earlier pre-0.1.40.2 run wedged fully (empty HTTP reply, needed a kill). Freezes only ever occurred during the prompt pass, never decode.
With STRATA_IO_PREFETCH=1 a full bench completed with no stalls (one data point — the prefetch path may be bypassing the fault storm).
- File tier re-reads tens of GB per request (STRATA_IO_STATS)
file tier I/O this request: the OS read 34,889.1 MB from storage (255,655 major faults)
for 89,694.7 MB of expert reads
file tier I/O this request: the OS read 24,306.9 MB ... page cache held 0.0 MB at the read
io prefetch: 917 read ahead (1329.3 MB by pread), 0 used, 0 unused
page cache held 0.0 MB even though ~17-20 GB is reclaimable — layer evictions (FADV_DONTNEED) make every request cold; 562 GB cumulative read for ~40 requests. Also 0 used / 0 unused for read-ahead looks like consumed preads aren't accounted for.
- Prefill degrades within a run (same binary, same config)
bench_prefill.py, back-to-back, engine otherwise idle:
| trial | prompt | prefill | decode |
|---|---|---|---|
| 1 | 4,210 | 75.5 tok/s | 16.9 |
| 2 | 8,830 | 87.5 tok/s | 21.0 |
| 3 | 4,210 | 22.2 | 5.5 |
| 4 | 8,830 | 33.5 | 34.9 |
Directional, not random: first two prompts fast, last two 3-4x slower. Follow-ups inherit it — 114-tok follow-up after a 4,210 fresh = 56-59 tok/s, after an 8,830 fresh = 15-17 tok/s. A large fresh prompt seems to leave the file tier re-streaming what it just evicted.
Ruled out
OOM (none in journalctl), NVMe errors (SMART pass, 1.5 GB/s reads), GPU Xid/HSA errors (none in dmesg), VRAM exhaustion (~890 MiB free at load), hipBLASLt table mismatch (confirmed loaded each start).
Notes
- ReBAR aperture is 256 MB (board DSDT zeroes the windows; local issue — mentioned only because the mmap fault path is what a small aperture penalizes)
- card0's memory clock sits at 96 MHz while DPM reads high mid-request (card1 stays at 1258 MHz); it does climb under full load
Questions
- When the prompt pass stalls at lock_mm_and_find_vma, is the watchdog catching the fault storm itself or a real hang?
- Is page cache held 0.0 MB expected with --mmap-experts + FADV_DONTNEED evictions?
- Are consumed prefetch preads miscounted (0 used, 0 unused)?
Happy to run A/B cells o
Activity
serve: the #29 watchdog waits (up to 10x, STRATA_WATCHDOG_IO_S) while…
fix the escape in the watchdog IO note (#1407)
Niko1221 commentedon Oct 8, 2026
Thanks, this is a useful report. What we can say:
STRATA_WATCHDOG_IO_S, 0 turns it off) and prints a line saying so. A real hang still stops at the old limit.For 32 GB of RAM with a 39 GB mmap pack, your own finding is the cheapest fix:
STRATA_IO_PREFETCH=1. We have not added a startup recommendation for it. Could you retest on 0.1.41 withSTRATA_IO_STATS=1, with and withoutSTRATA_IO_PREFETCH=1, 3+ interleaved runs of the 4,210 / 8,830-token trials?Iokkeh2025 commentedon Oct 8, 2026
A/B test results: STRATA_IO_PREFETCH on/off, v0.1.41, 32 GB RAM
Tested on v0.1.41 with STRATA_IO_STATS=1, 3 interleaved runs of 4,210 / 8,830-token prompts, engine otherwise idle. All 18 runs completed with zero crashes.
Conclusion: IO_PREFETCH provides NO measurable benefit on this system. The on/off numbers are statistically identical across all 3 runs.
What actually fixed the stalls: the 0.1.41 watchdog change. The old watchdog killed the engine during slow prompt reads (fault storm on the mmap file tier). The new watchdog watches major-fault count and bytes read, and waits up to 10x the limit while they grow. With that fix, the same hardware that was crashing every 4-8 requests now completes 18 consecutive runs without a single crash.
Recommendation: do not add a startup recommendation for IO_PREFETCH on 32 GB systems — it does not help. The watchdog fix in 0.1.41 is the real solution.
Testing note for users: close all other applications before benchmarking Strata. The best-case scenario is to SSH into the Strata server machine from a phone or laptop, leaving the server machine fully dedicated to the engine. Running a desktop session, browser, or other GPU/RAM-consuming apps alongside Strata will materially affect results.
Merge Niko1221/Strata main (v0.1.41) into main
Vop3p commentedon Oct 9, 2026
On the stall half of this (not the RAM re-read half, which I have no data on): the same symptom on a 2x RX 6900 XT (gfx1030) layer split - the #29 watchdog firing while reading a long prompt, expert pool idle - traced to the GPU's GFX power gating interacting with the adaptive expert swaps' SDMA copies (#884). Two things that made it go away there, both cheap to try on gfx1200:
STRATA_HIP_ADAPT_KERNEL_COPY=1(0.1.41): the swaps copy with a kernel instead of SDMA. Stock 0.1.41 split, 128K prompts: 0 stalls with it, no measurable speed cost (bench: community report, 2x RX 6900 XT (gfx1030) + Ryzen 5 5600X, IQ3_S at 131K on 0.1.41 - layer split: stock, the two gfx103x switches, the open PRs, and the decode window #1657,split-stock-kc)./sys/kernel/debug/dri/N/amdgpu_gfxoff(a reference count; read it back, 0 means off): 10/10 stalls to 0/10 on this box, same speed.If either one changes your stall rate, that would say whether the gfx1200 stall is the same animal.
Developed with an AI coding assistant; every number above was measured on 2x RX 6900 XT (gfx1030, PCIe 4.0 x8 each) / Ryzen 5 5600X, ROCm 10.0.
docs: the 0.1.41 watchdog, body-cap and batch-groups behaviors no pag…
docs: the 0.1.41 watchdog, body-cap and batch-groups behaviors no pag…
serve: ファイルの段を OS が読んでいる間は、#29 の見張りが最大 10 倍まで待つ（STRATA_WATCHDOG_IO_S、N…
Watchdog: a file tier busy within the last 10 s keeps the I/O allowan…
Niko1221 commentedon Oct 13, 2026
Thanks for the report. To keep the tracker current we are closing issues older than three days: 0.1.44 is out with a long list of fixes, and many older reports are covered by it or by 0.1.43.1.
If this still happens on 0.1.44, please leave a comment here with your engine log (and the version line it prints), and we will reopen it. A new issue that links this one works too.
見張り: ファイルの段が直前の 10 秒以内に読まれていれば読み込みの猶予を続ける（スワップに追われるホストには読みのない秒がある、Nik…
serve: the Niko1221#29 watchdog waits (up to 10x, STRATA_WATCHDOG_IO_…
fix the escape in the watchdog IO note (Niko1221#1407)
Watchdog: a file tier busy within the last 10 s keeps the I/O allowan…
metal: a batch slot window beats the serve watchdog, as on CUDA
ウォッチドッグ: 層が進んでいる間は止めず、止めるときは経った時間と理由を出す（#44）