From 0ffcfc7580385fc79a290a2fe80a7950086b5740 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 17:48:39 +0200 Subject: [PATCH 1/4] POCL: mark OpenCL calls as GC-safe Calls into PoCL can block for a long time: building programs, and waiting for kernels, which the POCL back-end does on every launch. As plain ccalls, they prevent every other thread from collecting garbage in the meantime. Use GPUToolbox's `@gcsafe_ccall` for all of them, as CUDA.jl and OpenCL.jl do. --- Project.toml | 2 ++ src/pocl/nanoOpenCL.jl | 54 ++++++++++++++++++++++-------------------- 2 files changed, 30 insertions(+), 26 deletions(-) diff --git a/Project.toml b/Project.toml index b7ee65576..755c70097 100644 --- a/Project.toml +++ b/Project.toml @@ -10,6 +10,7 @@ projects = ["test", "lib/KernelInterface", "lib/KernelInterface/test"] Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e" Atomix = "a9b6321e-bd34-4604-b9c9-b65b8de01458" GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55" +GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc" KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" LLVM = "929cbde3-209d-540e-8aea-75f648917ca0" MacroTools = "1914dd2f-81c6-5fcd-8719-6d5c9610ff09" @@ -42,6 +43,7 @@ Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.2.1" EnzymeCore = "0.7, 0.8.1" GPUCompiler = "2.7" +GPUToolbox = "3.1" KernelInterface = "0.4" LLVM = "9.9" LinearAlgebra = "1.6" diff --git a/src/pocl/nanoOpenCL.jl b/src/pocl/nanoOpenCL.jl index cec2e0deb..d7a67e949 100644 --- a/src/pocl/nanoOpenCL.jl +++ b/src/pocl/nanoOpenCL.jl @@ -4,6 +4,8 @@ using ..POCL: platform, device, context, queue import pocl_standalone_jll +using GPUToolbox: @gcsafe_ccall + using Printf const libopencl = pocl_standalone_jll.libpocl @@ -511,7 +513,7 @@ const cl_command_queue_properties = cl_bitfield const cl_event_info = cl_uint @checked function clGetPlatformIDs(num_entries, platforms, num_platforms) - @ccall libopencl.POclGetPlatformIDs( + @gcsafe_ccall libopencl.POclGetPlatformIDs( num_entries::cl_uint, platforms::Ptr{cl_platform_id}, num_platforms::Ptr{cl_uint} )::cl_int @@ -521,7 +523,7 @@ end platform, param_name, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetPlatformInfo( + @gcsafe_ccall libopencl.POclGetPlatformInfo( platform::cl_platform_id, param_name::cl_platform_info, param_value_size::Csize_t, param_value::Ptr{Cvoid}, @@ -530,7 +532,7 @@ end end @checked function clGetDeviceIDs(platform, device_type, num_entries, devices, num_devices) - @ccall libopencl.POclGetDeviceIDs( + @gcsafe_ccall libopencl.POclGetDeviceIDs( platform::cl_platform_id, device_type::cl_device_type, num_entries::cl_uint, devices::Ptr{cl_device_id}, num_devices::Ptr{cl_uint} @@ -541,7 +543,7 @@ end device, param_name, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetDeviceInfo( + @gcsafe_ccall libopencl.POclGetDeviceInfo( device::cl_device_id, param_name::cl_device_info, param_value_size::Csize_t, param_value::Ptr{Cvoid}, param_value_size_ret::Ptr{Csize_t} @@ -552,7 +554,7 @@ function clCreateContext( properties, num_devices, devices, pfn_notify, user_data, errcode_ret ) - return @ccall libopencl.POclCreateContext( + return @gcsafe_ccall libopencl.POclCreateContext( properties::Ptr{cl_context_properties}, num_devices::cl_uint, devices::Ptr{cl_device_id}, pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid}, @@ -561,11 +563,11 @@ function clCreateContext( end @checked function clReleaseContext(context) - @ccall libopencl.POclReleaseContext(context::cl_context)::cl_int + @gcsafe_ccall libopencl.POclReleaseContext(context::cl_context)::cl_int end function clCreateProgramWithIL(context, il, length, errcode_ret) - return @ccall libopencl.POclCreateProgramWithIL( + return @gcsafe_ccall libopencl.POclCreateProgramWithIL( context::cl_context, il::Ptr{Cvoid}, length::Csize_t, errcode_ret::Ptr{cl_int} @@ -573,14 +575,14 @@ function clCreateProgramWithIL(context, il, length, errcode_ret) end @checked function clReleaseProgram(program) - @ccall libopencl.POclReleaseProgram(program::cl_program)::cl_int + @gcsafe_ccall libopencl.POclReleaseProgram(program::cl_program)::cl_int end @checked function clBuildProgram( program, num_devices, device_list, options, pfn_notify, user_data ) - @ccall libopencl.POclBuildProgram( + @gcsafe_ccall libopencl.POclBuildProgram( program::cl_program, num_devices::cl_uint, device_list::Ptr{cl_device_id}, options::Ptr{Cchar}, pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid} @@ -591,7 +593,7 @@ end program, param_name, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetProgramInfo( + @gcsafe_ccall libopencl.POclGetProgramInfo( program::cl_program, param_name::cl_program_info, param_value_size::Csize_t, param_value::Ptr{Cvoid}, param_value_size_ret::Ptr{Csize_t} @@ -602,7 +604,7 @@ end program, device, param_name, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetProgramBuildInfo( + @gcsafe_ccall libopencl.POclGetProgramBuildInfo( program::cl_program, device::cl_device_id, param_name::cl_program_build_info, param_value_size::Csize_t, @@ -612,25 +614,25 @@ end end function clCreateKernel(program, kernel_name, errcode_ret) - return @ccall libopencl.POclCreateKernel( + return @gcsafe_ccall libopencl.POclCreateKernel( program::cl_program, kernel_name::Ptr{Cchar}, errcode_ret::Ptr{cl_int} )::cl_kernel end @checked function clReleaseKernel(kernel) - @ccall libopencl.POclReleaseKernel(kernel::cl_kernel)::cl_int + @gcsafe_ccall libopencl.POclReleaseKernel(kernel::cl_kernel)::cl_int end @checked function clSetKernelArg(kernel, arg_index, arg_size, arg_value) - @ccall libopencl.POclSetKernelArg( + @gcsafe_ccall libopencl.POclSetKernelArg( kernel::cl_kernel, arg_index::cl_uint, arg_size::Csize_t, arg_value::Ptr{Cvoid} )::cl_int end @checked function clSetKernelArgSVMPointer(kernel, arg_index, arg_value) - @ccall libopencl.POclSetKernelArgSVMPointer( + @gcsafe_ccall libopencl.POclSetKernelArgSVMPointer( kernel::cl_kernel, arg_index::cl_uint, arg_value::Ptr{Cvoid} )::cl_int @@ -640,7 +642,7 @@ end kernel, device, param_name, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetKernelWorkGroupInfo( + @gcsafe_ccall libopencl.POclGetKernelWorkGroupInfo( kernel::cl_kernel, device::cl_device_id, param_name::cl_kernel_work_group_info, param_value_size::Csize_t, @@ -654,7 +656,7 @@ end input_value, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetKernelSubGroupInfo( + @gcsafe_ccall libopencl.POclGetKernelSubGroupInfo( kernel::cl_kernel, device::cl_device_id, param_name::cl_kernel_sub_group_info, input_value_size::Csize_t, @@ -670,7 +672,7 @@ end command_queue, kernel, work_dim, global_work_size::NTuple{3, Csize_t}, local_work_size::NTuple{3, Csize_t}, event::Ref{cl_event} ) - @ccall libopencl.POclEnqueueNDRangeKernel( + @gcsafe_ccall libopencl.POclEnqueueNDRangeKernel( command_queue::cl_command_queue, kernel::cl_kernel, work_dim::cl_uint, C_NULL::Ptr{Csize_t}, @@ -688,7 +690,7 @@ end local_work_size, num_events_in_wait_list, event_wait_list, event ) - @ccall libopencl.POclEnqueueNDRangeKernel( + @gcsafe_ccall libopencl.POclEnqueueNDRangeKernel( command_queue::cl_command_queue, kernel::cl_kernel, work_dim::cl_uint, global_work_offset::Ptr{Csize_t}, @@ -701,7 +703,7 @@ end end function clCreateCommandQueue(context, device, properties, errcode_ret) - return @ccall libopencl.POclCreateCommandQueue( + return @gcsafe_ccall libopencl.POclCreateCommandQueue( context::cl_context, device::cl_device_id, properties::cl_command_queue_properties, errcode_ret::Ptr{cl_int} @@ -709,22 +711,22 @@ function clCreateCommandQueue(context, device, properties, errcode_ret) end @checked function clReleaseCommandQueue(command_queue) - @ccall libopencl.POclReleaseCommandQueue(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.POclReleaseCommandQueue(command_queue::cl_command_queue)::cl_int end @checked function clFinish(command_queue) - @ccall libopencl.POclFinish(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.POclFinish(command_queue::cl_command_queue)::cl_int end @checked function clWaitForEvents(num_events, event_list) - @ccall libopencl.POclWaitForEvents(num_events::cl_uint, event_list::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.POclWaitForEvents(num_events::cl_uint, event_list::Ptr{cl_event})::cl_int end @checked function clGetEventInfo( event, param_name, param_value_size, param_value, param_value_size_ret ) - @ccall libopencl.POclGetEventInfo( + @gcsafe_ccall libopencl.POclGetEventInfo( event::cl_event, param_name::cl_event_info, param_value_size::Csize_t, param_value::Ptr{Cvoid}, param_value_size_ret::Ptr{Csize_t} @@ -732,7 +734,7 @@ end end @checked function clReleaseEvent(event) - @ccall libopencl.POclReleaseEvent(event::cl_event)::cl_int + @gcsafe_ccall libopencl.POclReleaseEvent(event::cl_event)::cl_int end # Init @@ -1272,7 +1274,7 @@ end function set_arg!(k::Kernel, idx::Integer, arg::T) where {T} # `Ref{T}` makes `ccall` pass a pointer to a stack copy of `arg` - err = @ccall libopencl.POclSetKernelArg( + err = @gcsafe_ccall libopencl.POclSetKernelArg( k::cl_kernel, cl_uint(idx - 1)::cl_uint, sizeof(T)::Csize_t, arg::Ref{T} )::cl_int if err == CL_INVALID_ARG_SIZE From 6ad5ab3b743af0525297a97e1a1dbec03620b6ab Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 18:51:51 +0200 Subject: [PATCH 2/4] POCL: wait for kernels cooperatively MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit KernelAbstractions requires `synchronize` to be cooperative: waiting for the device should yield to other tasks rather than block the thread. The POCL back-end launches synchronously, because its arrays are plain `Array`s, but every launch waited for its kernel with a blocking `clWaitForEvents`, so no other task could run on the launching thread in the meantime. Wait using GPUToolbox's `cooperative_wait` instead: poll the kernel's status for at most 10 µs, and then have PoCL notify us of its completion through `clSetEventCallback`. Waiting on a worker thread, or polling for longer, would compete with the kernel for the CPU cores it executes on, making it considerably slower. The wait cannot be interrupted, as the kernel may be using memory the caller would release: interrupts are only thrown once the kernel completed. The first launch after new code has been defined allocates to look up the completion callback again, so measure allocations from a warmed-up function. --- Project.toml | 2 +- src/pocl/backend.jl | 12 ++++++--- src/pocl/nanoOpenCL.jl | 58 +++++++++++++++++++++++++++++++++++++++--- test/runtests.jl | 32 ++++++++++++++++++++++- 4 files changed, 95 insertions(+), 9 deletions(-) diff --git a/Project.toml b/Project.toml index 755c70097..a6c83be53 100644 --- a/Project.toml +++ b/Project.toml @@ -43,7 +43,7 @@ Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.2.1" EnzymeCore = "0.7, 0.8.1" GPUCompiler = "2.7" -GPUToolbox = "3.1" +GPUToolbox = "3.3" KernelInterface = "0.4" LLVM = "9.9" LinearAlgebra = "1.6" diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl index b2f2c6603..286e6b4df 100644 --- a/src/pocl/backend.jl +++ b/src/pocl/backend.jl @@ -154,13 +154,17 @@ end function KI.launch(obj::KI.Kernel{POCLBackend}, groups::Dims{3}, items::Dims{3}, args::Tuple) # the kernel only gets pointers to the arrays in `args` and captured by `f`, so keep # them alive until it completes. POCL launches synchronously, see the implementation - # note on `synchronize` + # note on `synchronize`. waiting for the kernel yields to other tasks, as `synchronize` + # should (see the documentation on its semantics). f = obj.kern.f - event = GC.@preserve f args begin + GC.@preserve f args begin event = POCL.launch_tuple(obj.kern.kernel, args; local_size = items, global_size = groups .* items) - wait(event) + try + wait(event) + finally + cl.clReleaseEvent(event) + end end - cl.clReleaseEvent(event) return nothing end diff --git a/src/pocl/nanoOpenCL.jl b/src/pocl/nanoOpenCL.jl index d7a67e949..257bd7b91 100644 --- a/src/pocl/nanoOpenCL.jl +++ b/src/pocl/nanoOpenCL.jl @@ -4,7 +4,7 @@ using ..POCL: platform, device, context, queue import pocl_standalone_jll -using GPUToolbox: @gcsafe_ccall +using GPUToolbox: GPUToolbox, @gcsafe_ccall, cooperative_wait using Printf @@ -718,6 +718,17 @@ end @gcsafe_ccall libopencl.POclFinish(command_queue::cl_command_queue)::cl_int end +@checked function clFlush(command_queue) + @gcsafe_ccall libopencl.POclFlush(command_queue::cl_command_queue)::cl_int +end + +@checked function clSetEventCallback(event, command_exec_callback_type, pfn_notify, user_data) + @gcsafe_ccall libopencl.POclSetEventCallback( + event::cl_event, command_exec_callback_type::cl_int, + pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid} + )::cl_int +end + @checked function clWaitForEvents(num_events, event_list) @gcsafe_ccall libopencl.POclWaitForEvents(num_events::cl_uint, event_list::Ptr{cl_event})::cl_int end @@ -1543,6 +1554,7 @@ struct Event end Base.unsafe_convert(::Type{cl_event}, e::Event) = e.id +const CL_EVENT_COMMAND_QUEUE = 0x11d0 const CL_EVENT_COMMAND_EXECUTION_STATUS = 0x11d3 function Base.getproperty(evt::Event, s::Symbol) @@ -1557,11 +1569,51 @@ function Base.getproperty(evt::Event, s::Symbol) end end +const CL_COMPLETE = 0 const CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST = -14 +# driver notification that a command has completed +function notify_completion(::cl_event, ::Cint, payload::Ptr{Cvoid}) + GPUToolbox.signal_completion(payload) + return +end +function subscribe_completion(evt, payload) + callback = @cfunction(notify_completion, Cvoid, (cl_event, Cint, Ptr{Cvoid})) + return clSetEventCallback(evt, CL_COMPLETE, callback, payload) +end + +# commands have completed when their execution status is `CL_COMPLETE`, or negative when +# they were terminated abnormally +iscomplete(evt::Event) = evt.status <= CL_COMPLETE + +blocking_wait(evt::Event) = unchecked_clWaitForEvents(cl_uint(1), Ref(evt.id)) + function Base.wait(evt::Event) - evt_id = Ref(evt.id) - err = unchecked_clWaitForEvents(cl_uint(1), evt_id) + # wait without blocking the thread, so that other tasks can run in the meantime. after + # polling briefly, PoCL notifies us when the command completes: waking a worker thread + # to wait for it, or polling for longer, would compete with the command for the CPU + # cores it executes on. + # + # this cannot be interrupted, as kernels may be using memory that callers would release: + # an interrupt is only thrown once the kernel has completed (or waiting failed, in which + # case we block), but host memory still needs to be synchronized before unwinding. + try + # commands only need to start executing once their queue has been flushed + queue = Ref{cl_command_queue}() + clGetEventInfo(evt, CL_EVENT_COMMAND_QUEUE, sizeof(cl_command_queue), queue, C_NULL) + clFlush(queue[]) + + cooperative_wait( + blocking_wait, evt; subscribe = subscribe_completion, isdone = iscomplete, + spin = 10.0e-6 + ) + catch + blocking_wait(evt) + rethrow() + end + + # synchronize host memory and report errors (without blocking anymore) + err = unchecked_clWaitForEvents(cl_uint(1), Ref(evt.id)) if err == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST error("Kernel execution failed") elseif err != CL_SUCCESS diff --git a/test/runtests.jl b/test/runtests.jl index 3beb38cbf..a0957b8bf 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -95,7 +95,11 @@ end mod.launch_few(few, A) mod.launch_many(many, A) @test all(==(sum(1:40)), A) - @test @allocated(mod.launch_many(many, A)) <= @allocated(mod.launch_few(few, A)) + # the first launch after new code has been defined (like the function below) allocates, + # to look up the callback that notifies it of the kernel's completion again + allocated(launch, k, A) = @allocated launch(k, A) + allocated(mod.launch_many, many, A) + @test allocated(mod.launch_many, many, A) <= allocated(mod.launch_few, few, A) end @testset "POCL compilation cache" begin @@ -160,6 +164,32 @@ end @test all(==(2.0f0), A) end +@kernel function busy_kernel!(A, n) + I = @index(Global) + acc = 0.0f0 + for j in 1:n + acc += sin(Float32(j) + acc) + end + @inbounds A[I] = acc +end + +# POCL launches wait for the kernel to finish, but let other tasks run in the meantime +@testset "POCL cooperative launches" begin + A = zeros(Float32, 1024) + kernel = busy_kernel!(POCLBackend()) + kernel(A, 1; ndrange = length(A)) # compile + n = 1000 + while @elapsed(kernel(A, n; ndrange = length(A))) < 0.1 + n *= 2 + end + + ran = Ref(false) + task = @async ran[] = true + kernel(A, n; ndrange = length(A)) + @test ran[] + wait(task) +end + # not part of the shared testsuite: not every back-end supports bits-union arrays @testset "POCL zeros/ones of bits-union types" begin for T in (Union{Missing, Bool}, Union{Missing, Int32}) From 05731167bd950256f4fcace9a3e52a0dd187a2cc Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 13:31:49 +0200 Subject: [PATCH 3/4] POCL: let the many-arguments allocation test wait by blocking Whether a launch allocates now depends on how it waits: one that polling finds completed doesn't, one that waits for the completion callback does. Which one a launch takes depends on how long the kernel runs, so the test comparing a launch with many arguments to one with few failed randomly on CI (e.g. 144 > 32 bytes). Add an internal switch to wait for kernels by blocking instead, and use it in that test, so that it only compares how the arguments are passed. --- src/pocl/nanoOpenCL.jl | 35 +++++++++++++++++++++-------------- test/runtests.jl | 13 +++++++++---- 2 files changed, 30 insertions(+), 18 deletions(-) diff --git a/src/pocl/nanoOpenCL.jl b/src/pocl/nanoOpenCL.jl index 257bd7b91..cd5288776 100644 --- a/src/pocl/nanoOpenCL.jl +++ b/src/pocl/nanoOpenCL.jl @@ -1588,6 +1588,10 @@ iscomplete(evt::Event) = evt.status <= CL_COMPLETE blocking_wait(evt::Event) = unchecked_clWaitForEvents(cl_uint(1), Ref(evt.id)) +# block the thread while waiting for commands instead, e.g., to make tests independent of +# how a wait was performed +const blocking_waits = Ref(false) + function Base.wait(evt::Event) # wait without blocking the thread, so that other tasks can run in the meantime. after # polling briefly, PoCL notifies us when the command completes: waking a worker thread @@ -1597,22 +1601,25 @@ function Base.wait(evt::Event) # this cannot be interrupted, as kernels may be using memory that callers would release: # an interrupt is only thrown once the kernel has completed (or waiting failed, in which # case we block), but host memory still needs to be synchronized before unwinding. - try - # commands only need to start executing once their queue has been flushed - queue = Ref{cl_command_queue}() - clGetEventInfo(evt, CL_EVENT_COMMAND_QUEUE, sizeof(cl_command_queue), queue, C_NULL) - clFlush(queue[]) - - cooperative_wait( - blocking_wait, evt; subscribe = subscribe_completion, isdone = iscomplete, - spin = 10.0e-6 - ) - catch - blocking_wait(evt) - rethrow() + if !blocking_waits[] + try + # commands only need to start executing once their queue has been flushed + queue = Ref{cl_command_queue}() + clGetEventInfo(evt, CL_EVENT_COMMAND_QUEUE, sizeof(cl_command_queue), queue, C_NULL) + clFlush(queue[]) + + cooperative_wait( + blocking_wait, evt; subscribe = subscribe_completion, isdone = iscomplete, + spin = 10.0e-6 + ) + catch + blocking_wait(evt) + rethrow() + end end - # synchronize host memory and report errors (without blocking anymore) + # synchronize host memory and report errors (without blocking anymore, unless waiting + # by blocking) err = unchecked_clWaitForEvents(cl_uint(1), Ref(evt.id)) if err == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST error("Kernel execution failed") diff --git a/test/runtests.jl b/test/runtests.jl index a0957b8bf..9411a9583 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -95,11 +95,16 @@ end mod.launch_few(few, A) mod.launch_many(many, A) @test all(==(sum(1:40)), A) - # the first launch after new code has been defined (like the function below) allocates, - # to look up the callback that notifies it of the kernel's completion again + # waiting for a kernel allocates when it involves a completion callback, which depends + # on how long the kernel takes, so measure launches that wait by blocking instead allocated(launch, k, A) = @allocated launch(k, A) - allocated(mod.launch_many, many, A) - @test allocated(mod.launch_many, many, A) <= allocated(mod.launch_few, few, A) + POCL.cl.blocking_waits[] = true + try + allocated(mod.launch_many, many, A) + @test allocated(mod.launch_many, many, A) <= allocated(mod.launch_few, few, A) + finally + POCL.cl.blocking_waits[] = false + end end @testset "POCL compilation cache" begin From c9697db66a0a6fe7e5f77ef194d9e01440dead34 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 15:03:39 +0200 Subject: [PATCH 4/4] Require GPUToolbox 3.3.2 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 3.3.2 keeps the task woken by a completion callback from spinning on a lock held by PoCL's callback thread after preempting it. With PoCL using every core, that took a core away from the next kernel, making 50-400 µs kernels 15-30% slower on the benchmark bot (JuliaGPU/GPUToolbox.jl#27). --- Project.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index a6c83be53..bf1c71b9f 100644 --- a/Project.toml +++ b/Project.toml @@ -43,7 +43,7 @@ Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.2.1" EnzymeCore = "0.7, 0.8.1" GPUCompiler = "2.7" -GPUToolbox = "3.3" +GPUToolbox = "3.3.2" KernelInterface = "0.4" LLVM = "9.9" LinearAlgebra = "1.6"