Skip to content

[Engine] HIP, two gfx1201 cards, layer split: batched prompt read stops at a chunk boundary, stage-2 card busy until reset; on this host triggered by a -50 mV offset, clean at -25 mV and 0 mV #4229

Closed
Closed
[Engine] HIP, two gfx1201 cards, layer split: batched prompt read stops at a chunk boundary, stage-2 card busy until reset; on this host triggered by a -50 mV offset, clean at -25 mV and 0 mV#4229

Description

@eysie101
eysie101
opened on Oct 26, 2026

Version 1.3, 2026-10-26

On a two-card HIP layer split, 11 of 23 fresh prompt reads of several chunks stopped at a chunk boundary; 0 of 2 on one card; 0 of 158 reads within one chunk (Q5_K_M, Q4_K_S).

Related: #541, #579, #697, #3297, #3443, #3680.

Environment

Item Value
Host HP Z8 G4, Xeon Gold 6134, Proxmox VE 9.2.21, kernel 7.0.14-25-pve
Guest Arch Linux, kernel 7.2.6-arch2-1 with its in-tree amdgpu, q35 11.0, cpu: host, 16 vCPU, 155648 MiB, both cards hostpci: pcie=1,driver=vfio
Cards 2x Radeon AI PRO R9700 32 GB, gfx1201; guest 01:00.0 (ASUS 1043:0626), 02:00.0 (Gigabyte 1458:242f)
Firmware, both cards VBIOS 115-G287BP00-100; MES 0x93, MES_KIQ 0x93, MEC 0xd7a, PFP 0xc76, ME 0xc12, RLC 0xbe7da0, SMC 104.80.0, SDMA0 0x798e96
Link guest reports 32.0 GT/s x16; engine PCIe probe 13.8 GB/s host to device
Card policy power1_cap 210 W, OD_VDDGFX_OFFSET -50 mV, power_dpm_force_performance_level manual; amdgpu: Overdrive is enabled
ROCm 7.2.4, hip-runtime-amd 7.2.4-1; engine libraries _rocm_sdk_libraries_gfx120X_all
Engine Strata 0.1.44.1 @296b59d, HIP build; 0.1.50 (source 6270e2d) in 2 runs
Model Qwen3.8-Flash-Next-Uncensored Q5_K_M, GGUF in 3 shards with its strata-pack; Q4_K_S and Q6_K where stated
Split layer_split: auto: layers 0 to 23 on device 0, 24 to 47 and the head on device 1
Client serve/server.py, /v1/chat/completions, parallel: 2, one request at a time unless stated

Engine arguments (Q5)

Argument Value
--pack / --native / --ple-gguf the strata-pack; GGUF shard 1 of 3; the same shard
--expert-profile data/expert-profile.bin
--max-context 262144
--kv / --kv-resident int8 / 32768
--expert-cache auto
--vram-reserve-mib / --vram-reserve-later-mib 3072 / 3072
--prefill auto (8192), 16384, 32768 as stated
--pool-workers 8
--resident-budget-gib 104
--ple-io direct
--spec / --spec-min-p 2 / 0
--suffix-draft / --lookup-chain / --mtp-max-t 0 / 0 / 0
--adapt-every / --adapt-async 0 / 0
--mtp / --batch-mtp / --batch-groups rt / set / 1
--pipeline-windows 0
env STRATA_ROUTE_TAIL_SKIP=0, STRATA_RESIDENT_HEADROOM_GIB=12

Engine start lines (Q5, two cards)

Line Value
PCIe probe 13.8 GB/s, pcie_frac 0.38 (default 0.55)
layer split auto K=24; the caches hold 13,318 of 24,576 profiled pairs; CUDA1 holds its weights, session [24, 48) and the head, 26.22 GiB free
pool 8 workers on CPUs 1 to 8, the host thread on CPU 0, aux threads on CPUs 9 to 15
hit path R4 hit path ON - resident experts are computed on the GPU
chunk 8192 96-slot ring; the prompt path borrows 1,065 CUDA0 slots and 1,074 of CUDA1's 6,097 (3.73 GiB each)
chunk 32768 96-slot ring; 3,635 and 3,662 slots (12.70 GiB each)

Observation

  • The server's token counter stops at a multiple of the chunk size (24,576; 32,768; 57,344) and does not move again.
  • The engine's stage string at the stop: reading the prompt (batched): waiting for the GPU (attention, router) at layer L of the prompt chunk from token N; L = 26 (N = 57,344, no watchdog, 852 s) and L = 40 (N = 24,576, watchdog); both layers lie in stage 2's range 24 to 47.
  • In the two 0.1.50 runs the public Progress object showed finishing the chunk from token N (N = 24,576 and 40,960).
  • The stage-2 card reads gpu_busy_percent 100, 74 to 106 W, sclk 3554 to 3577 MHz; its VRAM use does not change.
  • The stage-1 card reads 3%, 14 W.
  • The host main thread sits in hipStreamSynchronize; one host thread in kfd_wait_on_events; three in sched_yield.
  • The stall report: 0 layers served since the last finished step, expert pool: epoch 6892, batch epoch 6892: 27 of 27 jobs claimed, 27 done; 8 of 8 workers parked, 8 sleeping, host idle for 158097 ms, threads waiting on the disk (state D): 0 of 97.
  • With STRATA_WATCHDOG_S=120 the engine exits -6: no progress for 120 s during a request; without it the run does not end.
  • One run logged HW Exception by GPU node-1 (Agent handle: ...) reason :GPU Hang before the exit.
  • Teardown: MES(0) failed to respond to msg=REMOVE_QUEUE, msg=SUSPEND, failed to suspend gangs from MES, MES might be in unrecoverable state, issue a GPU reset, Failed to evict queue 8, GPU mode1 reset, VRAM is lost due to GPU reset!, GPU reset(N) succeeded!, [drm] device wedged, but no recovery needed.
  • The guest journal holds no GPUVM or page fault line for the whole boot (7 resets).
  • The host journal holds no IOMMU, AER or pcieport line between 2026-10-24 and 2026-10-26.
  • Twice after a teardown the card kept gpu_busy_percent 100, 80 W, 3570 MHz with no process holding /dev/kfd or /dev/dri/renderD*; cat /sys/kernel/debug/dri/0000:02:00.0/amdgpu_gpu_recover ended it.
  • rocgdb -p <pid> (rocm-gdb 7.2.4) on the stopped engine: fatal error: failed to suspend os_queue_id 9, amd-dbgapi: could not attach to process; guest kernel: amdgpu: PID <pid> not debug enabled for op 6; the card reset within 2 s.

Occurrence

Cards Chunk Prompt tokens Boundaries Completed
2 8192 65,536 7 2 of 7 (0.1.44.1), 0 of 1 (0.1.50)
2 8192, STRATA_PF_STEP_SYNC=1 65,536 7 2 of 2 (0.1.44.1), 0 of 1 (0.1.50)
2 8192, HIP_VISIBLE_DEVICES=1,0 65,536 7 0 of 1; the ASUS card is stage 2 and spins
2 8192, stages serialised (patch, see below) 65,536 7 0 of 1; stop at token 24,576
2 8192 16,384 1 1 of 1
2 8192, Q6_K 28,107 3 0 of 1
2 32768 27,182 to 30,966 0 5 of 5
2 32768 62,999 to 65,571 1 4 of 5; stop at token 32,768
2 32768, Q4_K_S 64,011 1 1 of 1
2 16384, Q6_K 30,412 and 64,011 1 and 3 2 of 2
2 16384, Q6_K 5,091, turn 4 of 30, prefix cached 0 0 of 1
2 16384, Q6_K up to 4,608 and session turns 0 65 of 66
2 8192 and 32768, Q5_K_M and Q4_K_S up to 4,608; 30-turn sessions; 2 concurrent 0 158 of 158
1 (gpu absent) 8192 65,536 7 2 of 2

No effect on the stop

  • STRATA_PF_STEP_SYNC=1: 1 stop in 3 runs.
  • HSA_USERPTR_FOR_PAGED_MEM=0 with GPU_PINNED_MIN_XFER_SIZE=1048576: 1 stop in 1 run.
  • Engine 0.1.50: 2 stops in 2 runs.
  • HIP_VISIBLE_DEVICES=1,0: the stop follows the stage-2 role.
  • Stages serialised by a source patch (stage 1 waits for stage 2's chunk before reading the next chunk; src/prefill/prefill.cpp @296b59d after hand_buf_ ^= 1;, gated by STRATA_PF_STAGE_SERIAL=1): 1 stop in 1 run.

Source (0.1.44.1 @296b59d)

  • doorbell_wait_kernel and copy_rows_from_mapped_after spin on host memory; src/prefill/prefill.cpp calls neither.
  • src/prefill/prefill.cpp uses the CPU pool only with STRATA_PREFILL_CPU_SHARE; it was not set.

Separate observation

Logs

engine-log-q5-L1-chunk8192.txt (last 80 of 129 lines) 
strata generate: layer split, CUDA1: 24.28 GiB free of 31.86, room for experts 21.19 GiB
strata generate: layer split: CUDA1 runs layers 24-47, expert cache 6097 slots (21.18 GiB), 6097 of its 12288 profiled pairs; slot 0 verified
strata generate: layer split: CUDA0 runs layers 0-23
strata generate: layer split auto: L2 probe: the split is startable on the live free figures
strata generate: routing-aware prefetch on (the next layer's router, file tier)
strata generate: R4 hit path ON - resident experts are computed on the GPU
strata generate: 8 expert-pool workers + the host thread
strata generate: session is up (engine 0.1.44.1)
strata generate: token graph hit path: 12608 resident experts, decided on the device
FileExpertSource: RAM budget 104.00 GiB: 11968 of the 41.68 GiB of experts the GPU cache does not hold, by profile rank; the rest are read from the files
FileExpertSource: allocating 49.14 GiB page-locked cache complement
FileExpertSource: 40.74 GiB really free for a 49.14 GiB page-locked complement; the model files' cached pages were released: 129.52 GiB free
FileExpertSource: copied cache complement through layer 8/48 (9.68 GiB, 7 s)
FileExpertSource: copied cache complement through layer 16/48 (18.19 GiB, 11 s)
FileExpertSource: copied cache complement through layer 24/48 (25.11 GiB, 16 s)
FileExpertSource: copied cache complement through layer 32/48 (32.16 GiB, 20 s)
FileExpertSource: copied cache complement through layer 40/48 (41.61 GiB, 26 s)
FileExpertSource: copied cache complement through layer 48/48 (49.14 GiB, 29 s)
FileExpertSource: mapped pinned cache complement ready: resident 49.14 GiB, pinned 49.14 GiB in 42 s; page-locked and mapped
FileExpertSource: the split stages lend 2139 slots to the prompt path: 2139 keep their experts in RAM too (7.45 GiB)
FileExpertSource: 6097 verified additional-GPU experts remain on the mmap fallback
strata generate: resident RAM mode: 49.14 GiB of experts in RAM (page-locked), 6511 in the GPU cache; adaptive swaps exchange them with the GPU cache (no file reads)
strata generate: the file tier reads through the file cache (re-checked with the RAM copy built (49.14 GiB): 83.1 GiB available, 0.0 GiB of it still to be taken by the RAM copy, 36.4 GiB of experts read from the files: the file cache can keep the ones that come back)
strata serve: prompt chunk auto: 8192 tokens, a 96-slot ring
strata serve: the prompt path borrows 1065 CUDA0 cache slots (3.73 GiB)
strata serve:   CUDA1 prompt path borrows 1074 of its 6097 slots (3.73 GiB)
strata serve: lend sizing: the prompt path borrows 2139 slots (ring 96 slots at chunk 8192), 2139 of them keep their experts in RAM too
strata hc: CUDA1: the hyper-connection read runs as staged (the norm per token and stream, the down projection's activations staged ahead by cp.async); checked bit for bit against the plain read on this card (STRATA_HC_SPLIT=1 or 0 for the earlier ones)
strata verify: window up to 8 tokens, 107.9 MiB of device buffers
strata serve: layer split: layers 0-23 (CUDA0), 24-47 (CUDA1), one hand-off per window
strata hc: CUDA0: the hyper-connection read runs as staged (the norm per token and stream, the down projection's activations staged ahead by cp.async); checked bit for bit against the plain read on this card (STRATA_HC_SPLIT=1 or 0 for the earlier ones)
strata verify: window up to 8 tokens, 107.9 MiB of device buffers
strata mtp: draft head over 106299 tokens (212.9 MiB)
strata verify: batch windows of up to 2 sequences (layers [0, 24))
strata verify: batch windows of up to 2 sequences (layers [24, 48))
strata serve: adaptive overlap: off: not on a layer split
strata serve: adaptive tier asynchronous (--adapt-async 1): up to 96 swaps per round, a round every 4 windows
strata serve: 2936 MiB of VRAM free with everything loaded
strata serve: CUDA1: 2808 MiB of VRAM free with everything loaded
strata verify: capturing the 2-token window (2936 MiB of VRAM free)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: capturing the 2-token window (2808 MiB of VRAM free)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: capturing the 1-token window (2898 MiB of VRAM free)
strata verify: captured the 1-token window (upload no error, sync no error)
strata verify: capturing the 1-token window (2774 MiB of VRAM free)
strata verify: captured the 1-token window (upload no error, sync no error)
strata decode timing: 16 windows, avg T 1.94, 1.81 tokens/window, 55.82 ms/window = verify 53.18 (GPU-reach wait 18.73 + per-layer host 6.80 [plan 0.18 actq 0.19 jobs 0.04 CPU 13.47] + stage 0.18) + commit/emit 0.73 + draft 1.48; per layer-window: CPU experts 1.97 (2.36 entries), VRAM hits 16.44, PCIe 0.55; adapt wait 0.42 join 0.00, 18.0 swaps/window
strata serve: prompt 59 tokens = 0 reused + 59 read in 1449 ms (40.7 tok/s), 29 generated in 893 ms (32.5 tok/s), drafts accepted 14 of 15, 1 checkpoints
strata serve: decode expert cache hit rate: 87.4% (12625 hits / 14440 lookups); 440 more read by the GPU over PCIe or from another GPU (3.0% of all 14880 routed)
strata serve: resident RAM: 49.14 GiB of experts in RAM, 226 exchanged with the VRAM tier, 0 blob reads from the file
strata serve: asynchronous adaptive tier: 4 rounds, 288 experts swapped in, 156.6 ms per round (start to flip, between windows)
strata serve: expert tiers: GPU 12625 hits this request; since the start RAM 5619 blobs, files 0 blobs 47106.0 MB read (the GGUF in place)
strata serve: file tier I/O this request: the OS read 94956.1 MB from storage (17465 major faults) for 47106.0 MB of expert reads
strata serve: KV streaming: 97.87% of 6198 block reads hit VRAM, 0.5 MiB read from RAM
strata decode timing: 36 windows, avg T 1.97, 1.78 tokens/window, 43.60 ms/window = verify 40.54 (GPU-reach wait 17.04 + per-layer host 2.79 [plan 0.15 actq 0.09 jobs 0.02 CPU 4.67] + stage 0.02) + commit/emit 0.74 + draft 1.52; per layer-window: CPU experts 0.82 (0.89 entries), VRAM hits 18.72, PCIe 0.11; adapt wait 0.80 join 0.00, 7.9 swaps/window
strata serve: prompt 2086 tokens = 0 reused + 2086 read in 7562 ms (275.9 tok/s), 64 generated in 1570 ms (40.8 tok/s), drafts accepted 28 of 35, 1 checkpoints
strata serve: decode expert cache hit rate: 95.5% (32344 hits / 33882 lookups); 198 more read by the GPU over PCIe or from another GPU (0.6% of all 34080 routed)
strata serve: resident RAM: 49.14 GiB of experts in RAM, 520 exchanged with the VRAM tier, 0 blob reads from the file
strata serve: asynchronous adaptive tier: 13 rounds, 668 experts swapped in, 146.5 ms per round (start to flip, between windows)
strata serve: expert tiers: GPU 32344 hits this request; since the start RAM 21108 blobs, files 0 blobs 47106.0 MB read (the GGUF in place)
strata serve: file tier I/O this request: the OS read 164.2 MB from storage (0 major faults) for 0.0 MB of expert reads
strata serve: KV streaming: 98.60% of 231060 block reads hit VRAM, 13.0 MiB read from RAM
strata decode timing: 77 windows, avg T 1.99, 1.66 tokens/window, 46.70 ms/window = verify 43.47 (GPU-reach wait 18.05 + per-layer host 3.34 [plan 0.14 actq 0.11 jobs 0.02 CPU 6.15] + stage 0.03) + commit/emit 0.70 + draft 1.79; per layer-window: CPU experts 1.07 (1.14 entries), VRAM hits 18.51, PCIe 0.21; adapt wait 0.74 join 0.00, 9.9 swaps/window
strata serve: prompt 16420 tokens = 0 reused + 16420 read in 18658 ms (880.0 tok/s), 128 generated in 3596 ms (35.6 tok/s), drafts accepted 52 of 76, 2 checkpoints
strata serve: decode expert cache hit rate: 94.2% (68430 hits / 72640 lookups); 800 more read by the GPU over PCIe or from another GPU (1.1% of all 73440 routed)
strata serve: resident RAM: 49.14 GiB of experts in RAM, 1120 exchanged with the VRAM tier, 0 blob reads from the file
strata serve: asynchronous adaptive tier: 32 rounds, 1486 experts swapped in, 145.4 ms per round (start to flip, between windows)
strata serve: expert tiers: GPU 68430 hits this request; since the start RAM 57862 blobs, files 0 blobs 47106.0 MB read (the GGUF in place)
strata serve: file tier I/O this request: the OS read 764.9 MB from storage (0 major faults) for 0.0 MB of expert reads
strata serve: KV streaming: 96.13% of 483702 block reads hit VRAM, 75.3 MiB read from RAM
strata decode timing: 78 windows, avg T 1.99, 1.64 tokens/window, 42.53 ms/window = verify 39.17 (GPU-reach wait 16.94 + per-layer host 1.41 [plan 0.13 actq 0.05 jobs 0.01 CPU 2.23] + stage 0.10) + commit/emit 0.67 + draft 2.08; per layer-window: CPU experts 0.40 (0.42 entries), VRAM hits 19.43, PCIe 0.03; adapt wait 0.61 join 0.00, 2.7 swaps/window
strata serve: prompt 65571 tokens = 0 reused + 65571 read in 54665 ms (1199.5 tok/s), 128 generated in 3317 ms (38.6 tok/s), drafts accepted 50 of 77, 5 checkpoints
strata serve: decode expert cache hit rate: 97.9% (72738 hits / 74300 lookups); 100 more read by the GPU over PCIe or from another GPU (0.1% of all 74400 routed)
strata serve: resident RAM: 49.14 GiB of experts in RAM, 1318 exchanged with the VRAM tier, 0 blob reads from the file
strata serve: asynchronous adaptive tier: 51 rounds, 1770 experts swapped in, 144.0 ms per round (start to flip, between windows)
strata serve: expert tiers: GPU 72738 hits this request; since the start RAM 175765 blobs, files 0 blobs 47106.0 MB read (the GGUF in place)
strata serve: file tier I/O this request: the OS read 1442.0 MB from storage (0 major faults) for 0.0 MB of expert reads
strata serve: KV streaming: 93.91% of 489900 block reads hit VRAM, 120.2 MiB read from RAM
HW Exception by GPU node-2 (Agent handle: 0x55d2ad5fe450) reason :GPU Hang
server-out-q5-L1-chunk8192.txt (last 40 of 94 lines) 
[strata] reading the prompt: 57,344 of 64,614 tokens, 77 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 87 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 97 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 107 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 117 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 127 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 137 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 147 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 157 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 167 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 177 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 187 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 197 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 207 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 217 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 227 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 237 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 247 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 257 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 267 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 277 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 287 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 297 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 307 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 317 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 327 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 337 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 347 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 357 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 367 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 377 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 387 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 397 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 407 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 417 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 427 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 437 s so far
[strata] reading the prompt: 57,344 of 64,614 tokens, 447 s so far
[strata] the engine stopped unexpectedly (exit code -6). The usual cause is running out of RAM: Linux then ends the biggest program (check: sudo dmesg | grep -i -E 'killed process|out of memory'); Windows slows down instead. Close other programs or use a smaller model (Q2_0 / IQ2_XS). The next request starts the engine again. Its log: /var/lib/claude-r3/q5-L1/engine.log
[strata] done: 0 tokens in 456 s (0.0 tok/s) (error, cancel=False)
kernel-log-q5-L1-teardown.txt (last 12 of 12 lines) 
2026-10-26T06:47:50+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU mode1 reset
2026-10-26T06:47:50+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU smu mode1 reset
2026-10-26T06:47:52+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU reset succeeded, trying to resume
2026-10-26T06:47:52+01:00 arch1 kernel: amdgpu 0000:02:00.0: VRAM is lost due to GPU reset!
2026-10-26T06:47:52+01:00 arch1 kernel: amdgpu 0000:02:00.0: program CP_MES_CNTL : 0x4000000
2026-10-26T06:47:52+01:00 arch1 kernel: amdgpu 0000:02:00.0: program CP_MES_CNTL : 0xc000000
2026-10-26T06:47:52+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU reset(3) succeeded!
2026-10-26T06:47:52+01:00 arch1 kernel: amdgpu 0000:02:00.0: [drm] device wedged, but no recovery needed
2026-10-26T06:48:05+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE
2026-10-26T06:48:05+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE
2026-10-26T06:48:05+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE
2026-10-26T06:48:05+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE
engine-log-q5-X1-stage-serial.txt (last 80 of 123 lines) 
strata generate: the GPU computes the experts in the cache; it rounds differently from the CPU,
                 so a reply can differ slightly from a run without the cache (same quality:
                 bench/results/2026-09-27-cache-parity).
                 PROFILE, ranked by routing frequency, no eviction.
strata generate: the profile fill asks for 256 pairs ahead (STRATA_FILL_AHEAD)
strata generate: pre-filled 6511 of 6511 slots from the profile in 13.6 s (1793 MB/s); slot 0 verified
strata generate: layer split, CUDA1: 24.28 GiB free of 31.86, room for experts 21.19 GiB
strata generate: layer split: CUDA1 runs layers 24-47, expert cache 6097 slots (21.18 GiB), 6097 of its 12288 profiled pairs; slot 0 verified
strata generate: layer split: CUDA0 runs layers 0-23
strata generate: layer split auto: L2 probe: the split is startable on the live free figures
strata generate: routing-aware prefetch on (the next layer's router, file tier)
strata generate: R4 hit path ON - resident experts are computed on the GPU
strata generate: 8 expert-pool workers + the host thread
strata generate: session is up (engine 0.1.44.1)
strata generate: token graph hit path: 12608 resident experts, decided on the device
FileExpertSource: RAM budget 104.00 GiB: 11968 of the 41.68 GiB of experts the GPU cache does not hold, by profile rank; the rest are read from the files
FileExpertSource: allocating 49.14 GiB page-locked cache complement
FileExpertSource: 43.86 GiB really free for a 49.14 GiB page-locked complement; the model files' cached pages were released: 120.49 GiB free
FileExpertSource: copied cache complement through layer 8/48 (8.88 GiB, 5 s)
FileExpertSource: copied cache complement through layer 16/48 (17.20 GiB, 9 s)
FileExpertSource: copied cache complement through layer 24/48 (24.71 GiB, 12 s)
FileExpertSource: copied cache complement through layer 32/48 (32.34 GiB, 16 s)
FileExpertSource: copied cache complement through layer 40/48 (40.20 GiB, 19 s)
FileExpertSource: copied cache complement through layer 48/48 (49.14 GiB, 22 s)
FileExpertSource: mapped pinned cache complement ready: resident 49.14 GiB, pinned 49.14 GiB in 37 s; page-locked and mapped
FileExpertSource: the split stages lend 2139 slots to the prompt path: 2139 keep their experts in RAM too (7.45 GiB)
FileExpertSource: 6097 verified additional-GPU experts remain on the mmap fallback
strata generate: resident RAM mode: 49.14 GiB of experts in RAM (page-locked), 6511 in the GPU cache; adaptive swaps off
strata generate: the file tier reads through the file cache (re-checked with the RAM copy built (49.14 GiB): 84.2 GiB available, 0.0 GiB of it still to be taken by the RAM copy, 36.4 GiB of experts read from the files: the file cache can keep the ones that come back)
strata serve: prompt chunk auto: 8192 tokens, a 96-slot ring
strata serve: the prompt path borrows 1065 CUDA0 cache slots (3.73 GiB)
strata serve:   CUDA1 prompt path borrows 1074 of its 6097 slots (3.73 GiB)
strata serve: lend sizing: the prompt path borrows 2139 slots (ring 96 slots at chunk 8192), 2139 of them keep their experts in RAM too
strata hc: CUDA1: the hyper-connection read runs as staged (the norm per token and stream, the down projection's activations staged ahead by cp.async); checked bit for bit against the plain read on this card (STRATA_HC_SPLIT=1 or 0 for the earlier ones)
strata verify: window up to 8 tokens, 107.9 MiB of device buffers
strata serve: layer split: layers 0-23 (CUDA0), 24-47 (CUDA1), one hand-off per window
strata hc: CUDA0: the hyper-connection read runs as staged (the norm per token and stream, the down projection's activations staged ahead by cp.async); checked bit for bit against the plain read on this card (STRATA_HC_SPLIT=1 or 0 for the earlier ones)
strata verify: window up to 8 tokens, 107.9 MiB of device buffers
strata mtp: draft head over 106299 tokens (212.9 MiB)
strata verify: batch windows of up to 2 sequences (layers [0, 24))
strata verify: batch windows of up to 2 sequences (layers [24, 48))
strata serve: adaptive overlap: off: not on a layer split
strata serve: 2936 MiB of VRAM free with everything loaded
strata serve: CUDA1: 2808 MiB of VRAM free with everything loaded
strata verify: capturing the 2-token window (2936 MiB of VRAM free)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: capturing the 2-token window (2808 MiB of VRAM free)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: capturing the 1-token window (2898 MiB of VRAM free)
strata verify: captured the 1-token window (upload no error, sync no error)
strata verify: capturing the 1-token window (2774 MiB of VRAM free)
strata verify: captured the 1-token window (upload no error, sync no error)
strata serve: prompt 59 tokens = 0 reused + 59 read in 1368 ms (43.1 tok/s), 32 generated in 748 ms (42.8 tok/s), drafts accepted 16 of 16, 1 checkpoints
strata serve: decode expert cache hit rate: 86.0% (13198 hits / 15344 lookups); 496 more read by the GPU over PCIe or from another GPU (3.1% of all 15840 routed)
strata serve: resident RAM: 49.14 GiB of experts in RAM, 0 exchanged with the VRAM tier, 0 blob reads from the file
strata serve: expert tiers: GPU 13198 hits this request; since the start RAM 5878 blobs, files 0 blobs 47106.0 MB read (the GGUF in place)
strata serve: file tier I/O this request: the OS read 86400.8 MB from storage (21103 major faults) for 47106.0 MB of expert reads
strata serve: KV streaming: 97.87% of 6486 block reads hit VRAM, 0.6 MiB read from RAM
strata serve: prompt 2086 tokens = 0 reused + 2086 read in 7591 ms (274.8 tok/s), 64 generated in 1577 ms (40.6 tok/s), drafts accepted 24 of 40, 1 checkpoints
strata serve: decode expert cache hit rate: 91.8% (35220 hits / 38375 lookups); 505 more read by the GPU over PCIe or from another GPU (1.3% of all 38880 routed)
strata serve: resident RAM: 49.14 GiB of experts in RAM, 0 exchanged with the VRAM tier, 0 blob reads from the file
strata serve: expert tiers: GPU 35220 hits this request; since the start RAM 23098 blobs, files 0 blobs 47106.0 MB read (the GGUF in place)
strata serve: file tier I/O this request: the OS read 118.9 MB from storage (0 major faults) for 0.0 MB of expert reads
strata serve: KV streaming: 98.77% of 261894 block reads hit VRAM, 13.0 MiB read from RAM
strata serve: no progress for 120 s during a request (reading the prompt (batched): waiting for the GPU (attention, router) at layer 40 of the prompt chunk from token 24576) - stopping the engine so the server starts it again (issue #29)
strata serve: stall report (engine 0.1.44.1): stage "reading the prompt (batched): waiting for the GPU (attention, router) at layer 40 of the prompt chunk from token 24576" for 120 s; 0 layers served since the last finished step (0 = stopped, more = slow)
  VRAM free when this request began: 2798 MiB
  expert pool: epoch 6892, batch epoch 6892: 27 of 27 jobs claimed, 27 done; 8 of 8 workers parked, 8 sleeping; mode 0
  expert pool threads: w0=sleeping w1=sleeping w2=sleeping w3=sleeping w4=sleeping w5=sleeping w6=sleeping w7=sleeping; host idle for 156096 ms
  verify window (last window, not the current stage): 2 tokens at position 2148, host at layer step 24; the GPU rang 24; flags: served 24, plan (A) 24, copies (B) 24
  threads waiting on the disk (state D): 0 of 97
  memory: 116612 MiB resident, 0 MiB in swap, 83062 MiB RAM available; 21103 major page faults so far
  2 s later:
  expert pool: epoch 6892, batch epoch 6892: 27 of 27 jobs claimed, 27 done; 8 of 8 workers parked, 8 sleeping; mode 0
  expert pool threads: w0=sleeping w1=sleeping w2=sleeping w3=sleeping w4=sleeping w5=sleeping w6=sleeping w7=sleeping; host idle for 158097 ms
  verify window (last window, not the current stage): 2 tokens at position 2148, host at layer step 24; the GPU rang 24; flags: served 24, plan (A) 24, copies (B) 24
  threads waiting on the disk (state D): 0 of 97
  memory: 116612 MiB resident, 0 MiB in swap, 83061 MiB RAM available; 21103 major page faults so far
strata: released the verify window's GPU waits (#267): the GPU finished in 0 ms
strata: released the verify window's GPU waits (#267): the GPU finished in 0 ms
server-out-q5-X1-stage-serial.txt (last 40 of 45 lines) 
         YOUR PC CAN BE SLOW OR STOP RESPONDING FOR 1-3 MINUTES NOW - this is normal.
         Please wait and don't close this window; the browser opens when it is ready.
[strata] filling the GPU's expert cache (6511 experts, 22.68 GiB of VRAM) ...
[strata] still starting (29 s) - please wait ...
[strata] almost ready ...
[strata] still starting (68 s) - please wait ...
[strata] thinking level: high (reasoning_effort in the config; a request, or the Chat settings shared with other apps, can set its own)
[strata] model aliases: local
ready: http://127.0.0.1:18190/v1  (OpenAI: /v1/chat/completions, Anthropic: /v1/messages, context 262144 tokens)
       open http://127.0.0.1:18190/ in a browser to chat; close this window to stop the model
[strata] reading the prompt: 54 of 59 tokens, 1 s so far
[strata] done: 32 tokens in 2 s (46.2 tok/s) (length, cancel=False), expert cache 86.0% hit (+3.1% of the routed experts over PCIe)
[strata] reading the prompt: 2,081 of 2,086 tokens, 7 s so far
[strata] thinking: 35 of max 64 tokens, 48.1 tok/s, 8 s
[strata] done: 64 tokens in 9 s (41.4 tok/s) (length, cancel=False), expert cache 91.8% hit (+1.3% of the routed experts over PCIe)
[strata] the reply reached max tokens while still thinking, so it has no answer: a thinking budget (reasoning_budget_tokens, in the request or in strata-<model>.json for every request) leaves room to answer
[strata] reading the prompt: 65,423 tokens, 10 s so far
[strata] reading the prompt: 16,384 of 65,423 tokens, 19 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 28 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 38 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 48 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 58 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 68 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 78 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 88 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 98 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 108 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 118 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 128 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 138 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 148 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 158 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 168 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 178 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 188 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 198 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 208 s so far
[strata] reading the prompt: 24,576 of 65,423 tokens, 218 s so far
[strata] the engine stopped unexpectedly (exit code -6). The engine stopped itself because it had stopped making progress - a hang it caught. Its log line: strata serve: no progress for 120 s during a request (reading the prompt (batched): waiting for the GPU (attention, router) at layer 40 of the prompt chunk from token 24576) - stopping the engine so the server starts it again (issue #29) - please report it at github.com/Niko1221/Strata/issues. The next request starts the engine again. Its log: /var/lib/claude-r3/q5-X1/engine.log
[strata] done: 0 tokens in 220 s (0.0 tok/s) (error, cancel=False)
kernel-log-q5-X1-teardown.txt (last 12 of 12 lines) 
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: MODE1 reset
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU mode1 reset
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU smu mode1 reset
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU reset succeeded, trying to resume
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: VRAM is lost due to GPU reset!
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: program CP_MES_CNTL : 0x4000000
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: program CP_MES_CNTL : 0xc000000
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: GPU reset(7) succeeded!
2026-10-26T13:32:31+01:00 arch1 kernel: amdgpu 0000:02:00.0: [drm] device wedged, but no recovery needed
2026-10-26T13:33:02+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE
2026-10-26T13:33:02+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE
2026-10-26T13:33:02+01:00 arch1 kernel: amdgpu 0000:02:00.0: MES(0) failed to respond to msg=REMOVE_QUEUE

Activity

  1. github-actions
    added
    needs-formatTitle or template does not follow CONTRIBUTING.md yet
    on Oct 26, 2026

  2. github-actions commented on Oct 26, 2026

    @github-actions
    github-actionsBot
    on Oct 26, 2026 – with GitHub Actions

    Thanks for the issue. To keep the tracker readable, every item needs an area prefix in the title and the parts of the issue form filled in. This one is missing:

    • Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other], then a short description. Example: "[Engine] decode stops after 2 minutes".

    Edit the issue (title and description) and this note goes away on its own. If it is still missing parts in 5 days it will be closed, which is not a rejection: fix it and comment /reopen.

    Details: CONTRIBUTING.md

  3. eysie101
    changed the title [-]HIP, two gfx1201 cards, layer split: the batched prompt read stops at a chunk boundary and the stage-2 card stays busy until a reset[/-] [+][Engine] HIP, two gfx1201 cards, layer split: the batched prompt read stops at a chunk boundary and the stage-2 card stays busy until a reset[/+] on Oct 26, 2026
  4. github-actions
    removed
    needs-formatTitle or template does not follow CONTRIBUTING.md yet
    on Oct 26, 2026

  5. github-actions commented on Oct 26, 2026

    @github-actions
    github-actionsBot
    on Oct 26, 2026 – with GitHub Actions

    Thanks, this follows the format now.

  6. eysie101 commented on Oct 26, 2026

    @eysie101
    eysie101
    on Oct 26, 2026
    Author

    Undervolt test, same host and engine as above, power1_cap 210 W on both cards throughout, --prefill auto (8192), production binary 0.1.44.1 @296b59d, four fresh 64K prompts per run (28 chunk boundaries):

    OD_VDDGFX_OFFSET runs 64K reads completed stops
    -50 mV the table above 12 of 23 multi-chunk reads 11
    0 mV 2 8 of 8 0
    -25 mV 5 (Q5 at chunk 8192 and 32768; Q6_K at 8192, 32768 and a 30-turn session) 12 of 12 0

    Each completed read: 1,185 to 1,750 tok/s; decode 40 to 46 tok/s. No GPU reset during these runs; the kernel log stays clean. The stage-serialisation patch above and the two env workarounds were not used in these runs.

    The MES / mode-1 reset signature appeared only at -50 mV. Probability that 56 boundaries pass at the -50 mV rate is below 1 %.

    At -25 mV the Q6_K cases that failed at -50 mV also passed: the 30-turn session (hang at turn 4 before; now 31 of 31, no non-finite retry) and chunk 32768 (two #879 non-finite failures before; now 13 of 13, decode 31 to 38 tok/s instead of 9 to 12).

  7. eysie101
    changed the title [-][Engine] HIP, two gfx1201 cards, layer split: the batched prompt read stops at a chunk boundary and the stage-2 card stays busy until a reset[/-] [+][Engine] HIP, two gfx1201 cards, layer split: batched prompt read stops at a chunk boundary, stage-2 card busy until reset; on this host triggered by a -50 mV offset, clean at -25 mV and 0 mV[/+] on Oct 26, 2026

  8. eysie101 commented on Oct 26, 2026

    @eysie101
    eysie101
    on Oct 26, 2026
    Author

    Closing from the reporter's side: on this host the stops were caused by the -50 mV OD_VDDGFX_OFFSET, not by the engine. With -25 mV (and 0 mV) the same two-card layer split completed every multi-chunk read, the 30-turn sessions and the Q6_K chunk-32768 case that produced non-finite logits before. The measurements stay above for anyone who undervolts a gfx1201 pair (see also #3680). If a note in the troubleshooting docs is wanted, the voltage table above is the source.

  9. eysie101
    closed this as completedon Oct 26, 2026

  10. sdroege commented on Oct 31, 2026

    @sdroege
    sdroege
    on Oct 31, 2026

    Single-card confirmation on a Radeon AI PRO R9700 (gfx1201), engine 0.1.54, UD-Q4_K_XL, no layer split.

    At OD_VDDGFX_OFFSET -60 mV, batched prompt reads of 16K and 32K stopped at a chunk boundary: no progress for 60 s ... waiting for the GPU (attention, router) at layer L of the prompt chunk from token N, then MES(0) failed to respond to msg=REMOVE_QUEUE, MODE1 reset, VRAM lost. 6 of 7 runs. Independent of --prefill chunk (8192/16384), --vram-reserve-mib (576/2048), the gfx12 speed switches and --spec (2/4/8).

    At -25 mV and 0 mV: two full runs each (16K and 32K, 3 requests per stage) with no hang, no non-finite logits and no GPU reset.

    Devcoredump of the hang: [gfxhub] Page fault observed, faulty address 0x0, GCVM protection fault status 0, MES ring Rptr 0x0 Wptr 0x3c800, gfx ring idle (CP_RB_RPTR = CP_RB_WPTR = 0x200). So this reproduces your -50 mV finding on a single card without a layer split, with a different model and engine version.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions