Repository navigation
Some finalizers wait for all GPU work #1167
Description
Several finalizers make HIP or library calls that wait for all GPU work on the device, not just the work using the object. In the ROCm runtime, ihipFree calls SyncAllStreams() for memory that doesn't come from a pool (hipamd/src/hip_memory.cpp), which covers hipHostFree, and hipHostUnregister does the same. That stalls whatever thread happens to collect garbage, and deadlocks when the GPU work waits for a task on that thread (#537 was this with hostcalls).
Measured on a gfx1036 iGPU (ROCm 7.2.4, Julia 1.12), with a ~1 s kernel running on the main task's stream and the release triggered from a GC on thread 2 (medians of 7; a full GC costs ~90–120 ms on its own):
| release | blocks the collecting thread | worst launch on another thread |
|---|---|---|
device ROCArray (pool, hipFreeAsync) |
≈ GC only | ≈ GC only |
host ROCArray (hipHostFree) |
902 ms | ≈ GC only |
unsafe_wrapped host memory |
≈ GC, but the deferred hipHostUnregister later blocks thread 1 for 817 ms |
– |
rocRAND RNG after use, rocSPARSE MatInfo |
902 ms | ≈ GC only |
rocblas_destroy_handle, miopenDestroy, rocfft_plan_destroy (on handle-cache eviction) |
902 ms | 0.1 ms |
| module, stream, event, rocBLAS handle returned to the cache | ≈ GC only | ≈ GC only |
With a kernel spinning on a flag that a task on the same thread sets after a GC, collecting a host ROCArray deadlocks 5/5 (until the kernel's 5 s cap). Collecting a wrapped array deadlocks intermittently, when the deferred unregister callback runs before the flag-setting task.
Unlike with CUDA (JuliaGPU/CUDA.jl#3318), launches on other threads aren't held up, and device memory is already fine. So only the releases that can wait would need to be deferred: queue them from the finalizer, and perform them on reclaim, on out-of-memory and on device_synchronize, possibly reusing freed host buffers by size so pinned memory doesn't grow in the meantime.
Activity