100% Single Thread Usage on Ryzen AI MAX 395 (gfx1151, Strix Halo)

[RandomLemon]

When the engine start to prefill/decode, one or two CPU core's usage increase to 100%.

top:

    PID USER      PR  NI    VIRT    RES    SHR S  %CPU  %MEM     TIME+ COMMAND                                                                              
    327 root      20   0  107.2g  94.3g 296092 S 200.0  75.5  34:25.52 luce_server

Here is what I caught with perf record -g -p <PID>:

  Children      Self  Command      Shared Object                      Symbol                                                                                
+   62.49%    62.35%  luce_server  libhsa-runtime64.so.1.21.0         [.] rocr::core::Runtime::AsyncEventsLoop(void*)
+   33.36%    33.27%  luce_server  libhsa-runtime64.so.1.21.0         [.] rocr::core::InterruptSignal::WaitRelaxed(hsa_signal_condition_t, long, unsigned long, hsa_wait_state_t)

Here is what I caught with gdb -p <PID> and bt:

#0  0x00007c1079ef2525 in rocr::core::InterruptSignal::WaitRelaxed(hsa_signal_condition_t, long, unsigned long, hsa_wait_state_t) ()
   from /opt/rocm/core-7.14//lib/libhsa-runtime64.so.1
#1  0x00007c1079ef25c7 in rocr::core::InterruptSignal::WaitAcquire(hsa_signal_condition_t, long, unsigned long, hsa_wait_state_t) ()
   from /opt/rocm/core-7.14//lib/libhsa-runtime64.so.1
#2  0x00007c1079ee2d44 in rocr::HSA::hsa_signal_wait_scacquire(hsa_signal_s, hsa_signal_condition_t, long, unsigned long, hsa_wait_state_t) ()
   from /opt/rocm/core-7.14//lib/libhsa-runtime64.so.1

Here is my environment:

Debian 13.7
Ryzen AI MAX+ 395, with 128G RAM
ROCm 7.14.1

Here is my parameters:

export MODEL_DIR=/mnt/data/ai/llm/deepseek-v4-flash
export ROCM_PATH=/opt/rocm/core-7.14/
export HIP_PATH=$ROCM_PATH
export LD_LIBRARY_PATH=$ROCM_PATH/lib:$LD_LIBRARY_PATH
DFLASH_DS4_SPEC=1 \
DFLASH_DS4_DRAFT=$MODEL_DIR/DeepSeek-V4-Flash-0731-DSpark-draft-Q4RMFP4-denseF16.gguf \
DFLASH_DS4_SPARSE_DECODE_FLASH=1 \
./server/build-hip/luce_server $MODEL_DIR/DeepSeek-V4-Flash-0731-ROCMFPX-MIX-STRIX.gguf \
  --port 8002 \
  --target-device hip:0 \
  --max-ctx 500000 \
  --chunk 8192 \
  --cache-type-k q4_0 --cache-type-v q4_0 \
  --ds4-fused-decode --ds4-fused-verify-f16-kv \
  --ds4-expert-top-k 6 \
  --ds4-prefill sparse

[Graffioh]

Thanks for the report.

i did a quick investigation:

Thread1 (WaitRelaxed)

it's the default AMD mechanism for locking with a spin in these cases during a synchronization call - there is another tipology that could be used, a block (sleep), but this introduce a bit of overhead across the sync calls and also probably some side effects (that i didn't honestly measure)

• AMD HIP docs, device flags (the main reference): https://rocm.docs.amd.com/projects/HIP/en/latest/reference/hip_runtime_api/modules/device_management.html. The hipDeviceScheduleSpin entry says: "HIP runtime will actively spin in the thread which submitted the work until the command completes. This offers the lowest latency, but will consume a CPU core and may increase power."
• HIP source, spin is the default:
  https://github.com/ROCm/rocm-systems/blob/820ea79c1848e1291204e7f7e56ec68bd049704e/projects/clr/hipamd/src/hip_context.cpp#L52 sets SetActiveWait(true) for every device at startup.
• llama.cpp precedent: CUDA: Changing the CUDA scheduling strategy to spin ggml-org/llama.cpp#16585. On NVIDIA's GB10 iGPU, the non-spin default cost ~15% throughput from sync wake-up delay, so llama.cpp forces spin mode there. Lucebox inherits that code in its vendored ggml.

i suggest to keep this as default.

Thread2 (AsyncEventsLoop)

it's a reported AMD ROCm bug, where there is an active wait instead of a sleep (so the opposite of thread 1 default mechanism lol) regarding post-GPU signal callbacks

• [Bug] gfx1151: bundled ROCR 1.21 AsyncEventsLoop busy-spins one full core permanently after any GPU op (nightly AND 7.14 production wheels) — release-channel ROCR 1.18 clean on same kernel ROCm/TheRock#7051
• ROCr AsyncEventsLoop threads busy-spin (~2 CPU cores) for the lifetime of any GPU context under WSL2 — reproduced on system ROCm 7.2.4; interrupt-backed signal events apparently unavailable through /dev/dxg ROCm/librocdxg#60
• here it is a user report signaling a fix using this setup for his particular use case, and i assume it is the ROCm nightly version: HostQueue::finish CPU-waits inside the execution lock (small-batch marker path); a cross-thread hipStreamWaitEvent on the same stream blocks for the whole wait and deadlocks host-dependent kernels ROCm/rocm-systems#11678

both these threads cores are at 100% during inference work, not during idle status, so if we don't see any real overhead that this introduces on CPU side, i'll say we are fine as is and if you manage to find a good fix for the Thread2 case feel free to update us!

[RandomLemon]

Thanks for the reply! The 100% CPU usage issue causes my machine noisy and hot (about 95°C). So I did more test:

7.2 from pytorch: TWO CPU core 100% usage, with the AsyncEventsLoop issue
7.2.4 from AMD build: same as above
7.14.1 from AMD build: same
10.0.0 from AMD build: same
7.13.0 from https://rocm.nightlies.amd.com/v2/gfx1151/rocm/: same
10.1.0-a20260822 from https://rocm.nightlies.amd.com/whl-multi-arch/: ONE CPU core 100% usage, occupied by rocr::core::InterruptSignal::WaitRelaxed

Only the Nightly build 10.1.0-a20260822 has no hot issue. The WaitRelaxed thread occupied the cpu core but the cpu is only 70 °C compared to the original 95°C. So this might be a ROCm issue, and we may wait for AMD to release the 10.1 version. The solution comes from this: Comfy-Org/ComfyUI#15347

My environment is Debian13 LXC inside a ProxmoxVE, with kernel 7.0.14 and 6.17.13. The kernel version didn't affect the result.