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

[seanthegeek]

Summary

Under WSL2, any process that touches the GPU through ROCm spins two rocr::core::Runtime::AsyncEventsLoop threads at ~91% CPU each (~182% of a core total) for the entire lifetime of its GPU context — completely idle, no work submitted. GPU compute itself is correct and fast; the cost is ~2 host CPU cores burned for as long as a context exists. Reproduced with stock system ROCm 7.2.4 and independently with the ROCm 7.2 runtime bundled by Ollama 0.30.7 — same behavior on both, so this is not a packaging issue.

Environment

• GPU: AMD Radeon RX 7900 XTX (24 GB, RDNA 3, gfx1100), Adrenalin 26.6.1
• Host: Windows 11 + WSL2, kernel 6.6.114.1-microsoft-standard-WSL2, Ubuntu 24.04
• librocdxg.so.1.2.0, libdxcore.so
• Reproduced against two HSA runtimes: system ROCm 7.2.4 (libhsa-runtime64.so.1.18.70204) and Ollama 0.30.7's bundled ROCm 7.2 (libhsa-runtime64.so.1.18.70201)

Minimal reproduction (20 lines, no dependencies beyond hipcc)

#include <hip/hip_runtime.h>
#include <cstdio>
#include <unistd.h>

__global__ void noop() {}

int main() {
    if (hipSetDevice(0) != hipSuccess) { fprintf(stderr, "no device\n"); return 1; }
    void* p = nullptr;
    if (hipMalloc(&p, 1 << 20) != hipSuccess) { fprintf(stderr, "alloc failed\n"); return 1; }
    noop<<<1, 1>>>();
    if (hipDeviceSynchronize() != hipSuccess) { fprintf(stderr, "sync failed\n"); return 1; }
    fprintf(stderr, "GPU touched; idling 120s, pid=%d\n", getpid());
    sleep(120);
    hipFree(p);
    return 0;
}

hipcc --offload-arch=gfx1100 -o hipidle hipidle.hip && ./hipidle
# then watch it with top -H while it sleeps

Measured while the program sits in sleep(120): 181.8% CPU — two threads at 90.9% each, wchan=0 (never blocked in the kernel). The spin starts at context creation and ends only when the process exits.

Attribution

The two hot threads are ROCr's async-event handlers. With Ollama's bundled runtime (which ships symbol names) the stacks are explicit — both spinners sit in AsyncEventsLoop, and every other thread in the process is properly parked in cond_wait/accept4:

Thread 24 (LWP 328349):
#0  rocr::core::Runtime::AsyncEventsLoop(void*) ()   from libhsa-runtime64.so.1
#1  rocr::os::ThreadTrampoline(void*) ()             from libhsa-runtime64.so.1

Thread 23 (LWP 328350):
#0  rocr::core::Runtime::AsyncEventsLoop(void*) ()   from libhsa-runtime64.so.1
#1  rocr::os::ThreadTrampoline(void*) ()

The stripped system runtime shows the identical signature on the minimal reproduction — two threads, two nearby offsets in libhsa-runtime64.so.1, each one frame above start_thread:

Thread 3 (LWP 336579):
#0  0x00007ffff569ac37 in ?? () from /opt/rocm-7.2.4/lib/libhsa-runtime64.so.1
#1  0x00007ffff56f9a31 in ?? () from /opt/rocm-7.2.4/lib/libhsa-runtime64.so.1
#2  start_thread (...) at ./nptl/pthread_create.c:447

Thread 2 (LWP 336578):
#0  0x00007ffff569ac84 in ?? () from /opt/rocm-7.2.4/lib/libhsa-runtime64.so.1
#1  0x00007ffff56f9a31 in ?? () from /opt/rocm-7.2.4/lib/libhsa-runtime64.so.1
#2  start_thread (...) at ./nptl/pthread_create.c:447

Neither spinner ever enters a kernel wait — whatever wait they call returns immediately, so the event loop degenerates to a userspace poll.

Supporting evidence from the original (Ollama) observation

The bug was first noticed on Ollama 0.30.7: a GGUF model merely loaded to GPU (zero requests) pins the same two threads; evicting the model returns CPU to baseline instantly. Reloading the same model with GPU offload disabled (num_gpu: 0) idles at 0.0% CPU — same binary, same server — so the spin is exclusively on the GPU path.

Full captured thread dump from the Ollama runner (main thread + first 80 lines of thread apply all bt 8 — includes both spinners and representative parked threads; lower-numbered threads fell outside the capture window, and per-thread CPU sampling showed no thread other than LWPs 328349/328350 above 1%)

--- main thread (tid 328346) ---
#0  0x0000706b99098d71 in __futex_abstimed_wait_common64 (private=0, cancel=true, abstime=0x7ffd2ad39a70, op=393, expected=0, futex_word=0x15761b98) at ./nptl/futex-internal.c:57
#1  __futex_abstimed_wait_common (cancel=true, private=0, abstime=0x7ffd2ad39a70, clockid=0, expected=0, futex_word=0x15761b98) at ./nptl/futex-internal.c:87
#2  __GI___futex_abstimed_wait_cancelable64 (futex_word=futex_word@entry=0x15761b98, expected=expected@entry=0, clockid=clockid@entry=0, abstime=abstime@entry=0x7ffd2ad39a70, private=private@entry=0) at ./nptl/futex-internal.c:139
#3  0x0000706b9909bc8e in __pthread_cond_wait_common (abstime=0x7ffd2ad39a70, clockid=0, mutex=0x15761b48, cond=0x15761b70) at ./nptl/pthread_cond_wait.c:503
#4  ___pthread_cond_timedwait64 (cond=0x15761b70, mutex=0x15761b48, abstime=0x7ffd2ad39a70) at ./nptl/pthread_cond_wait.c:652
#5  0x0000706b9a3ca9f1 in server_queue::start_loop(long) () from /usr/local/lib/ollama/libllama-server-impl.so
#6  0x0000706b9a2b1608 in llama_server(int, char**) () from /usr/local/lib/ollama/libllama-server-impl.so
#7  0x0000706b9902a1ca in __libc_start_call_main (...) at ../sysdeps/nptl/libc_start_call_main.h:58
#8  0x0000706b9902a28b in __libc_start_main_impl (...) at ../csu/libc-start.c:360
#9  0x000000000040107e in ?? ()

Thread 25 (Thread 0x706b98fff6c0 (LWP 328348) "llama-server"):
#0  0x0000706b99098d71 in __futex_abstimed_wait_common64 (private=0, cancel=true, abstime=0x0, op=393, expected=0, futex_word=0x14e4b0e8) at ./nptl/futex-internal.c:57
#1  __futex_abstimed_wait_common (cancel=true, private=0, abstime=0x0, clockid=0, expected=0, futex_word=0x14e4b0e8) at ./nptl/futex-internal.c:87
#2  __GI___futex_abstimed_wait_cancelable64 (futex_word=futex_word@entry=0x14e4b0e8, expected=expected@entry=0, clockid=clockid@entry=0, abstime=abstime@entry=0x0, private=private@entry=0) at ./nptl/futex-internal.c:139
#3  0x0000706b9909b7ed in __pthread_cond_wait_common (abstime=0x0, clockid=0, mutex=0x14e4b090, cond=0x14e4b0c0) at ./nptl/pthread_cond_wait.c:503
#4  ___pthread_cond_wait (cond=0x14e4b0c0, mutex=0x14e4b090) at ./nptl/pthread_cond_wait.c:627
#5  0x0000706b994b9047 in std::condition_variable::wait(std::unique_lock<std::mutex>&) () from /lib/x86_64-linux-gnu/libstdc++.so.6
#6  0x0000706b99e51233 in common_log::resume()::{lambda()#1}::operator()() const () from /usr/local/lib/ollama/libllama-common.so.0
#7  0x0000706b994ecdb4 in ?? () from /lib/x86_64-linux-gnu/libstdc++.so.6

Thread 24 (Thread 0x706b1a3ff6c0 (LWP 328349) "llama-server"):
#0  0x0000706b98286aa6 in rocr::core::Runtime::AsyncEventsLoop(void*) () from /usr/local/lib/ollama/rocm_v7_2/libhsa-runtime64.so.1
#1  0x0000706b982e34dd in rocr::os::ThreadTrampoline(void*) () from /usr/local/lib/ollama/rocm_v7_2/libhsa-runtime64.so.1
#2  0x0000706b9909caa4 in start_thread (arg=<optimized out>) at ./nptl/pthread_create.c:447
#3  0x0000706b99129c6c in clone3 () at ../sysdeps/unix/sysv/linux/x86_64/clone3.S:78

Thread 23 (Thread 0x706b19bfe6c0 (LWP 328350) "llama-server"):
#0  0x0000706b982868b7 in rocr::core::Runtime::AsyncEventsLoop(void*) () from /usr/local/lib/ollama/rocm_v7_2/libhsa-runtime64.so.1
#1  0x0000706b982e34dd in rocr::os::ThreadTrampoline(void*) () from /usr/local/lib/ollama/rocm_v7_2/libhsa-runtime64.so.1
#2  0x0000706b9909caa4 in start_thread (arg=<optimized out>) at ./nptl/pthread_create.c:447
#3  0x0000706b99129c6c in clone3 () at ../sysdeps/unix/sysv/linux/x86_64/clone3.S:78

Thread 22 (Thread 0x706b18bfc6c0 (LWP 328352) "llama-server"):
#0  0x0000706b9912b914 in accept4 (fd=7, addr=..., addr_len=0x0, flags=524288) at ../sysdeps/unix/sysv/linux/accept4.c:31
#1  0x0000706b9a3e5d71 in httplib::Server::listen_internal() () from /usr/local/lib/ollama/libllama-server-impl.so
#2  0x0000706b994ecdb4 in ?? () from /lib/x86_64-linux-gnu/libstdc++.so.6
#3  0x0000706b9909caa4 in start_thread (arg=<optimized out>) at ./nptl/pthread_create.c:447
#4  0x0000706b99129c6c in clone3 () at ../sysdeps/unix/sysv/linux/x86_64/clone3.S:78

Threads 21-17 (LWPs 328353-328357): all identical —
#0-#4 __futex_abstimed_wait_common64 / pthread_cond_wait
#5  std::condition_variable::wait
#6  httplib::ThreadPool::worker(bool) () from /usr/local/lib/ollama/libllama-server-impl.so

Suspected mechanism

On bare-metal Linux, AsyncEventsLoop blocks in the kernel via interrupt-backed KFD signal events. The polling fallback that pins a core when interrupt-backed events are unavailable is long-known (ROCm/ROCm#748). Under WSL2 there is no /dev/kfd; the thunk layer is librocdxg over /dev/dxg — and the observed behavior matches interrupt-backed signal events not being available on this path, leaving every signal on the memory-polling fallback. Notably, both reproducing runtimes postdate the ROCm 6.3.2 AsyncEventsLoop anti-spin improvements, which apparently cannot engage here.

Impact

~2 CPU cores burned continuously for the lifetime of any GPU context under WSL2. For inference servers this is severe: Ollama's default keep-alive holds a model (and the spin) for minutes after every request, and a server keeping a model warm spins indefinitely.

---

🤖 This investigation and report were generated with Claude Code; the measurements, bisect, and thread dumps were captured on the reporter's machine and reviewed by the reporter.

[seanthegeek]

Root cause turned out to be in the HSA runtime, not librocdxg: when librocdxg's hsaKmtCreateEvent stub leaves a signal without an interrupt-backed event, ROCr's AsyncEventsLoop (core/runtime/runtime.cpp) falls into a no-sleep polling loop and spins a core per async-events thread. Opened a fix against rocm-systems that naps with an escalating backoff in that polling branch — measured idle CPU 182% → 1.07% on this RX 7900 XTX / WSL2 setup with no throughput regression: ROCm/rocm-systems#7898

[fcui-amd]

ROCm/rocm-systems@96d0b30 . hope this would help with your issue

[seanthegeek]

Thanks @fcui-amd, I really appreciate you turning this around so fast! I built and tested the patch on the original hardware, and unfortunately it doesn't seem to resolve the spin on this configuration — wanted to share the results in case they're useful.

I built the rocm-7.2.4 tag with the patch applied (from the 7.2.4 tag specifically, so the ABI matches the installed librocdxg and the patch is the only variable) and re-ran the hipidle repro from the issue on the same machine — RX 7900 XTX / gfx1100, ROCm 7.2.4, WSL2. Whole-process CPU over a fixed 10s idle window:

Runtime	Idle CPU	Spinner threads
Stock 7.2.4	183.5%	91.2% / 91.1%
7.2.4 + this patch	183.2%	91.2% / 91.1%
7.2.4 + unconditional sleep in Signal::WaitMultiple	182.1%	91.2% / 91.1%
7.2.4 + sleep in AsyncEventsLoop polling fallback	16.0%	8.1% / 8.0%

(The patched uSleep call is confirmed present in the built binary's WaitMultiple disassembly — it just isn't reached on this path.)

Why I think it misses: the patch throttles the polling fallback in Signal::WaitMultiple, but Runtime::AsyncEventsLoop (the two hot threads in the original dump) only calls WaitMultiple when flag().wait_any() is true, and wait_any_ defaults to false — it's only set by HSA_WAIT_ANY_DEBUG=1, in both 7.2.4 and develop. In the default configuration AsyncEventsLoop instead takes its own PrepareInterrupt / WaitForInterrupt path. Under WSL the signals have no interrupt-backed events (EopEvent() == nullptr), so PrepareInterrupt returns false, polling is set, interrupt_wait stays false, and the inner while (!finish) loop spins with no KFD wait and no sleep.

The last two rows isolate this: an unconditional sleep inside WaitMultiple changes nothing (so that function isn't on the hot path here), whereas throttling the AsyncEventsLoop polling fallback brings idle CPU from ~183% to ~16%.

A minimal change that worked for me, in runtime.cpp right after the if (interrupt_wait) WaitForInterrupt(); block in AsyncEventsLoop:

} else if (polling && !finish) {
  // No interrupt-backed events (e.g. WSL/dxg): throttle the memory-polling
  // fallback, which would otherwise busy-spin a full core.
  os::uSleep(20);
}

20 µs gives ~16%; a larger interval lowers it further at the cost of a little wakeup latency. I already have an open PR for the runtime side of this — ROCm/rocm-systems#7898— which puts the nap in exactly this AsyncEventsLoop polling branch, but with an escalating backoff (20 µs doubling to a 2 ms cap, reset to the floor on every processed event) rather than a flat sleep. That idles at ~1% CPU and leaves back-to-back kernel+sync throughput unchanged (verified on the same hardware). Happy to run any other variant you'd like measured.

Would you be open to taking ROCm/rocm-systems#7898 instead of the WaitMultiple change here? And happy to adjust it however you'd prefer.

---

Disclaimer: this comment was written by Claude (Anthropic's Claude Code) at my direction. The build, the hipidle runs, and all CPU measurements above were performed on my machine and reviewed by me before posting.

[fcui-amd]

You're right, and your fix already covers this. I'll drop my change

[seanthegeek]

@fcui-amd Any channce of getting ROCm/rocm-systems#7898 reviewed and merged before the next release? ROCm on WSL is impractical without it.