Repetetive AMDGPU.ones calls crash runtime

[pxl-th]

Following code snippet will crash:

julia> using AMDGPU

julia> for _ in 1:100_000
           AMDGPU.ones(Float32, 16);
       end

with this error:

ERROR: schedule: Task not runnable
Stacktrace:
  [1] error(s::String)
    @ Base ./error.jl:35
  [2] schedule(t::Task, arg::Any; error::Bool)
    @ Base ./task.jl:829
  [3] schedule
    @ ./task.jl:827 [inlined]
  [4] uv_writecb_task(req::Ptr{Nothing}, status::Int32)
    @ Base ./stream.jl:1166
  [5] poptask(W::Base.IntrusiveLinkedListSynchronized{Task})
    @ Base ./task.jl:963
  [6] wait()
    @ Base ./task.jl:972
  [7] wait(c::Base.GenericCondition{ReentrantLock})
    @ Base ./condition.jl:124
  [8] wait(e::Base.Event)
    @ Base ./lock.jl:457
  [9] wait(signal::AMDGPU.Runtime.ROCKernelSignal)
    @ AMDGPU.Runtime ~/.julia/dev/AMDGPU/src/runtime/kernel-signal.jl:51
 [10] gpu_call(::AMDGPU.ROCArrayBackend, f::Function, args::Tuple{ROCVector{Float32}, Float32}, threads::Int64, blocks::Int64; name::Nothing)
    @ AMDGPU ~/.julia/dev/AMDGPU/src/array.jl:16
 [11] #gpu_call#1
    @ ~/.julia/packages/GPUArrays/fqD8z/src/device/execution.jl:65 [inlined]
 [12] gpu_call
    @ ~/.julia/packages/GPUArrays/fqD8z/src/device/execution.jl:34 [inlined]
 [13] fill!
    @ ~/.julia/packages/GPUArrays/fqD8z/src/host/construction.jl:14 [inlined]
 [14] ones
    @ ~/.julia/dev/AMDGPU/src/array.jl:405 [inlined]
 [15] top-level scope
    @ ./REPL[2]:2

Putting logs on loots of stuff I noticed that (at least) ROCSignal associated with respective ROCKernelSignals are not killed. Their finalizers do not run.
I suspect that there may be some dangling references that prevent finalizers from being called.

One thing to note, that the same code, but for CUDA runs fine on a much weaker GPU.

Only when exiting Julia REPL they are triggered. Manually calling GC.gc() also does not help.

julia> 
[ Info: [main] RT exiting true
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8d3ded4000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8d3ded8000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8d3deda000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8d3dee4000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8d3def0000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8d3def6000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8dbc7af000))
[ Info: [signal] kill
[ Info: [kersig] kill ROCSignal(AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f8dcd6fa000))
[ Info: [signal] kill
[ Info: [queue] kill

[pxl-th]

For some reason, after thread in the errormonitor finishes waiting it still prevents kernel signal from being GC'd...

One way to fix this is to wait not in the separate thread, but directly:

julia> using AMDGPU

julia> dummy_kern() = (return nothing;)
dummy_kern (generic function with 1 method)

julia> wait(@roc dummy_kern())
[ Info: [kern] exception size 4096 bytes
[ Info: [queue] create
[ Info: [queue] error monitor wait
true

julia> GC.gc()
[ Info: [kersig] kill
[ Info: [signal] kill AMDGPU.HSA.LibHSARuntime.hsa_signal_s(0x00007f60976f8000)

With this change the above example also works fine.
But this requires us to wait at some point to perform cleanup that is done in _wait function.

Another way is to pass to thread in errormonitor a WeakRef(kernel_signal), then we'd extract the value and do the regular waiting.
This allows GC to collect kernel signals, but with the above for-loop example it still fails, probably because it launches too many threads...
But in reality you probably won't launch this many threads all at once...

schedule: Task not runnable
    Stacktrace:
      [1] error(s::String)
        @ Base ./error.jl:35
      [2] schedule(t::Task, arg::Any; error::Bool)
        @ Base ./task.jl:829
      [3] schedule
        @ ./task.jl:827 [inlined]
      [4] uv_writecb_task(req::Ptr{Nothing}, status::Int32)
        @ Base ./stream.jl:1166
      [5] poptask(W::Base.IntrusiveLinkedListSynchronized{Task})
        @ Base ./task.jl:963
      [6] wait()
        @ Base ./task.jl:972
      [7] wait(c::Base.GenericCondition{Base.Threads.SpinLock})
        @ Base ./condition.jl:124
      [8] wait_readnb(x::Base.TTY, nb::Int64)
        @ Base ./stream.jl:416
      [9] eof(s::Base.TTY)
        @ Base ./stream.jl:106
     [10] eof(io::REPL.Terminals.TTYTerminal)
        @ Base ./io.jl:450

[pxl-th]

So here's MWE that shows that signals are not GC'd until Julia exits.
Observe that only B objects are GC'd after we empty! both arrays.

mutable struct A
    x::Int64
end

function A()
    a = A(1)
    errormonitor(Threads.@spawn println(a))
    finalizer(a) do ai
        println("Finalizer for $ai")
    end
    a
end

mutable struct B
    x::Int64
end

function B()
    b = B(1)
    finalizer(b) do bi
        println("Finalizer for $bi")
    end
    b
end

function main()
    as = A[]
    bs = B[]

    push!(as, A())
    push!(bs, B())
    sleep(1)
    GC.gc()

    empty!(as)
    empty!(bs)
    GC.gc()

    println("done")
end
main()

[pxl-th]

Looks like related to JuliaLang/julia#40626