From 2e2e46056f6beb0584512bd03c67a2aafd3311f6 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 21:23:32 +0200 Subject: [PATCH 1/7] Mark OpenCL API calls as GC-safe Many OpenCL calls can block for a long time: waiting for commands, compiling programs, blocking transfers. As plain ccalls, they prevent every other thread from collecting garbage in the meantime. They can even deadlock: if a driver thread enters Julia (e.g., to run a callback) while a Julia thread blocks in the driver waiting on it, the GC waits on the blocked thread, and the driver thread waits on the GC (see JuliaGPU/CUDA.jl#2261). As in CUDA.jl, make all API calls with GPUToolbox's `@gcsafe_ccall`, including extension functions, which are called through function pointers (supported since GPUToolbox 3.2). On Julia versions without native GC-safe calls, `@gcsafe_ccall` converts arguments itself before passing them to `ccall`, which converts them again. Make that second, identity conversion unambiguous for `PtrOrCLPtr` and `RefOrCLRef`, as CUDA.jl does. --- Project.toml | 2 +- lib/cl/api.jl | 5 +- lib/cl/libopencl.jl | 1040 +++++++++++++++++++++++-------------------- lib/cl/pointer.jl | 6 + res/wrap.jl | 7 +- 5 files changed, 573 insertions(+), 487 deletions(-) diff --git a/Project.toml b/Project.toml index 0236d89d..9ab45ad4 100644 --- a/Project.toml +++ b/Project.toml @@ -33,7 +33,7 @@ SPIRVIntrinsics = {path = "lib/intrinsics"} Adapt = "4" GPUArrays = "11.2.1" GPUCompiler = "2.9" -GPUToolbox = "3.1" +GPUToolbox = "3.2" KernelAbstractions = "0.9.38" LLVM = "9.6" LinearAlgebra = "1" diff --git a/lib/cl/api.jl b/lib/cl/api.jl index 533a1add..cd2d3288 100644 --- a/lib/cl/api.jl +++ b/lib/cl/api.jl @@ -1,4 +1,5 @@ import OpenCL_jll +using GPUToolbox: @gcsafe_ccall const libopencl = OpenCL_jll.libopencl @@ -89,7 +90,7 @@ function ext_ccall(platform, ex) @gensym fptr esc(quote $fptr = $clGetExtensionFunctionAddressForPlatform($platform, $fn) - @ccall $(Expr(:($), fptr))($(argexprs...))::$ret + @gcsafe_ccall $(Expr(:($), fptr))($(argexprs...))::$ret end) end @@ -183,7 +184,7 @@ const initialized = Ref{Bool}(false) withenv("OCL_ICD_FILENAMES"=>ocd_filenames) do num_platforms = Ref{Cuint}() - @ccall libopencl.clGetPlatformIDs( + @gcsafe_ccall libopencl.clGetPlatformIDs( 0::cl_uint, C_NULL::Ptr{cl_platform_id}, num_platforms::Ptr{cl_uint})::cl_int diff --git a/lib/cl/libopencl.jl b/lib/cl/libopencl.jl index d8a47b8b..46a8c69b 100644 --- a/lib/cl/libopencl.jl +++ b/lib/cl/libopencl.jl @@ -232,485 +232,526 @@ end const cl_name_version = _cl_name_version @checked function clGetPlatformIDs(num_entries, platforms, num_platforms) - @ccall libopencl.clGetPlatformIDs(num_entries::cl_uint, platforms::Ptr{cl_platform_id}, - num_platforms::Ptr{cl_uint})::cl_int + @gcsafe_ccall libopencl.clGetPlatformIDs(num_entries::cl_uint, + platforms::Ptr{cl_platform_id}, + num_platforms::Ptr{cl_uint})::cl_int end @checked function clGetPlatformInfo(platform, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetPlatformInfo(platform::cl_platform_id, - param_name::cl_platform_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetPlatformInfo(platform::cl_platform_id, + param_name::cl_platform_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetDeviceIDs(platform, device_type, num_entries, devices, num_devices) - @ccall libopencl.clGetDeviceIDs(platform::cl_platform_id, device_type::cl_device_type, - num_entries::cl_uint, devices::Ptr{cl_device_id}, - num_devices::Ptr{cl_uint})::cl_int + @gcsafe_ccall libopencl.clGetDeviceIDs(platform::cl_platform_id, + device_type::cl_device_type, + num_entries::cl_uint, devices::Ptr{cl_device_id}, + num_devices::Ptr{cl_uint})::cl_int end @checked function clGetDeviceInfo(device, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetDeviceInfo(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})::cl_int + @gcsafe_ccall libopencl.clGetDeviceInfo(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})::cl_int end @checked function clCreateSubDevices(in_device, properties, num_devices, out_devices, num_devices_ret) - @ccall libopencl.clCreateSubDevices(in_device::cl_device_id, - properties::Ptr{cl_device_partition_property}, - num_devices::cl_uint, - out_devices::Ptr{cl_device_id}, - num_devices_ret::Ptr{cl_uint})::cl_int + @gcsafe_ccall libopencl.clCreateSubDevices(in_device::cl_device_id, + properties::Ptr{cl_device_partition_property}, + num_devices::cl_uint, + out_devices::Ptr{cl_device_id}, + num_devices_ret::Ptr{cl_uint})::cl_int end @checked function clRetainDevice(device) - @ccall libopencl.clRetainDevice(device::cl_device_id)::cl_int + @gcsafe_ccall libopencl.clRetainDevice(device::cl_device_id)::cl_int end @checked function clReleaseDevice(device) - @ccall libopencl.clReleaseDevice(device::cl_device_id)::cl_int + @gcsafe_ccall libopencl.clReleaseDevice(device::cl_device_id)::cl_int end @checked function clSetDefaultDeviceCommandQueue(context, device, command_queue) - @ccall libopencl.clSetDefaultDeviceCommandQueue(context::cl_context, - device::cl_device_id, - command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.clSetDefaultDeviceCommandQueue(context::cl_context, + device::cl_device_id, + command_queue::cl_command_queue)::cl_int end @checked function clGetDeviceAndHostTimer(device, device_timestamp, host_timestamp) - @ccall libopencl.clGetDeviceAndHostTimer(device::cl_device_id, - device_timestamp::Ptr{cl_ulong}, - host_timestamp::Ptr{cl_ulong})::cl_int + @gcsafe_ccall libopencl.clGetDeviceAndHostTimer(device::cl_device_id, + device_timestamp::Ptr{cl_ulong}, + host_timestamp::Ptr{cl_ulong})::cl_int end @checked function clGetHostTimer(device, host_timestamp) - @ccall libopencl.clGetHostTimer(device::cl_device_id, - host_timestamp::Ptr{cl_ulong})::cl_int + @gcsafe_ccall libopencl.clGetHostTimer(device::cl_device_id, + host_timestamp::Ptr{cl_ulong})::cl_int end function clCreateContext(properties, num_devices, devices, pfn_notify, user_data, errcode_ret) - @ccall libopencl.clCreateContext(properties::Ptr{cl_context_properties}, - num_devices::cl_uint, devices::Ptr{cl_device_id}, - pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_context + @gcsafe_ccall libopencl.clCreateContext(properties::Ptr{cl_context_properties}, + num_devices::cl_uint, + devices::Ptr{cl_device_id}, + pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_context end function clCreateContextFromType(properties, device_type, pfn_notify, user_data, errcode_ret) - @ccall libopencl.clCreateContextFromType(properties::Ptr{cl_context_properties}, - device_type::cl_device_type, - pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_context + @gcsafe_ccall libopencl.clCreateContextFromType(properties::Ptr{cl_context_properties}, + device_type::cl_device_type, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_context end @checked function clRetainContext(context) - @ccall libopencl.clRetainContext(context::cl_context)::cl_int + @gcsafe_ccall libopencl.clRetainContext(context::cl_context)::cl_int end @checked function clReleaseContext(context) - @ccall libopencl.clReleaseContext(context::cl_context)::cl_int + @gcsafe_ccall libopencl.clReleaseContext(context::cl_context)::cl_int end @checked function clGetContextInfo(context, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetContextInfo(context::cl_context, param_name::cl_context_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetContextInfo(context::cl_context, + param_name::cl_context_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clSetContextDestructorCallback(context, pfn_notify, user_data) - @ccall libopencl.clSetContextDestructorCallback(context::cl_context, - pfn_notify::Ptr{Cvoid}, - user_data::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetContextDestructorCallback(context::cl_context, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid})::cl_int end function clCreateCommandQueueWithProperties(context, device, properties, errcode_ret) - @ccall libopencl.clCreateCommandQueueWithProperties(context::cl_context, - device::cl_device_id, - properties::Ptr{cl_queue_properties}, - errcode_ret::Ptr{cl_int})::cl_command_queue + @gcsafe_ccall libopencl.clCreateCommandQueueWithProperties(context::cl_context, + device::cl_device_id, + properties::Ptr{cl_queue_properties}, + errcode_ret::Ptr{cl_int})::cl_command_queue end @checked function clRetainCommandQueue(command_queue) - @ccall libopencl.clRetainCommandQueue(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.clRetainCommandQueue(command_queue::cl_command_queue)::cl_int end @checked function clReleaseCommandQueue(command_queue) - @ccall libopencl.clReleaseCommandQueue(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.clReleaseCommandQueue(command_queue::cl_command_queue)::cl_int end @checked function clGetCommandQueueInfo(command_queue, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetCommandQueueInfo(command_queue::cl_command_queue, - param_name::cl_command_queue_info, - param_value_size::Csize_t, - param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetCommandQueueInfo(command_queue::cl_command_queue, + param_name::cl_command_queue_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end function clCreateBuffer(context, flags, size, host_ptr, errcode_ret) - @ccall libopencl.clCreateBuffer(context::cl_context, flags::cl_mem_flags, size::Csize_t, - host_ptr::Ptr{Cvoid}, errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateBuffer(context::cl_context, flags::cl_mem_flags, + size::Csize_t, + host_ptr::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreateSubBuffer(buffer, flags, buffer_create_type, buffer_create_info, errcode_ret) - @ccall libopencl.clCreateSubBuffer(buffer::cl_mem, flags::cl_mem_flags, - buffer_create_type::cl_buffer_create_type, - buffer_create_info::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateSubBuffer(buffer::cl_mem, flags::cl_mem_flags, + buffer_create_type::cl_buffer_create_type, + buffer_create_info::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreateImage(context, flags, image_format, image_desc, host_ptr, errcode_ret) - @ccall libopencl.clCreateImage(context::cl_context, flags::cl_mem_flags, - image_format::Ptr{cl_image_format}, - image_desc::Ptr{cl_image_desc}, host_ptr::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateImage(context::cl_context, flags::cl_mem_flags, + image_format::Ptr{cl_image_format}, + image_desc::Ptr{cl_image_desc}, + host_ptr::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreatePipe(context, flags, pipe_packet_size, pipe_max_packets, properties, errcode_ret) - @ccall libopencl.clCreatePipe(context::cl_context, flags::cl_mem_flags, - pipe_packet_size::cl_uint, pipe_max_packets::cl_uint, - properties::Ptr{cl_pipe_properties}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreatePipe(context::cl_context, flags::cl_mem_flags, + pipe_packet_size::cl_uint, + pipe_max_packets::cl_uint, + properties::Ptr{cl_pipe_properties}, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreateBufferWithProperties(context, properties, flags, size, host_ptr, errcode_ret) - @ccall libopencl.clCreateBufferWithProperties(context::cl_context, - properties::Ptr{cl_mem_properties}, - flags::cl_mem_flags, size::Csize_t, - host_ptr::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateBufferWithProperties(context::cl_context, + properties::Ptr{cl_mem_properties}, + flags::cl_mem_flags, size::Csize_t, + host_ptr::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreateImageWithProperties(context, properties, flags, image_format, image_desc, host_ptr, errcode_ret) - @ccall libopencl.clCreateImageWithProperties(context::cl_context, - properties::Ptr{cl_mem_properties}, - flags::cl_mem_flags, - image_format::Ptr{cl_image_format}, - image_desc::Ptr{cl_image_desc}, - host_ptr::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateImageWithProperties(context::cl_context, + properties::Ptr{cl_mem_properties}, + flags::cl_mem_flags, + image_format::Ptr{cl_image_format}, + image_desc::Ptr{cl_image_desc}, + host_ptr::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end @checked function clRetainMemObject(memobj) - @ccall libopencl.clRetainMemObject(memobj::cl_mem)::cl_int + @gcsafe_ccall libopencl.clRetainMemObject(memobj::cl_mem)::cl_int end @checked function clReleaseMemObject(memobj) - @ccall libopencl.clReleaseMemObject(memobj::cl_mem)::cl_int + @gcsafe_ccall libopencl.clReleaseMemObject(memobj::cl_mem)::cl_int end @checked function clGetSupportedImageFormats(context, flags, image_type, num_entries, image_formats, num_image_formats) - @ccall libopencl.clGetSupportedImageFormats(context::cl_context, flags::cl_mem_flags, - image_type::cl_mem_object_type, - num_entries::cl_uint, - image_formats::Ptr{cl_image_format}, - num_image_formats::Ptr{cl_uint})::cl_int + @gcsafe_ccall libopencl.clGetSupportedImageFormats(context::cl_context, + flags::cl_mem_flags, + image_type::cl_mem_object_type, + num_entries::cl_uint, + image_formats::Ptr{cl_image_format}, + num_image_formats::Ptr{cl_uint})::cl_int end @checked function clGetMemObjectInfo(memobj, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetMemObjectInfo(memobj::cl_mem, param_name::cl_mem_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetMemObjectInfo(memobj::cl_mem, param_name::cl_mem_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetImageInfo(image, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetImageInfo(image::cl_mem, param_name::cl_image_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetImageInfo(image::cl_mem, param_name::cl_image_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetPipeInfo(pipe, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetPipeInfo(pipe::cl_mem, param_name::cl_pipe_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetPipeInfo(pipe::cl_mem, param_name::cl_pipe_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clSetMemObjectDestructorCallback(memobj, pfn_notify, user_data) - @ccall libopencl.clSetMemObjectDestructorCallback(memobj::cl_mem, - pfn_notify::Ptr{Cvoid}, - user_data::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetMemObjectDestructorCallback(memobj::cl_mem, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid})::cl_int end function clSVMAlloc(context, flags, size, alignment) - @ccall libopencl.clSVMAlloc(context::cl_context, flags::cl_svm_mem_flags, size::Csize_t, - alignment::cl_uint)::CLPtr{Cvoid} + @gcsafe_ccall libopencl.clSVMAlloc(context::cl_context, flags::cl_svm_mem_flags, + size::Csize_t, + alignment::cl_uint)::CLPtr{Cvoid} end function clSVMFree(context, svm_pointer) - @ccall libopencl.clSVMFree(context::cl_context, svm_pointer::PtrOrCLPtr{Cvoid})::Cvoid + @gcsafe_ccall libopencl.clSVMFree(context::cl_context, + svm_pointer::PtrOrCLPtr{Cvoid})::Cvoid end function clCreateSamplerWithProperties(context, sampler_properties, errcode_ret) - @ccall libopencl.clCreateSamplerWithProperties(context::cl_context, - sampler_properties::Ptr{cl_sampler_properties}, - errcode_ret::Ptr{cl_int})::cl_sampler + @gcsafe_ccall libopencl.clCreateSamplerWithProperties(context::cl_context, + sampler_properties::Ptr{cl_sampler_properties}, + errcode_ret::Ptr{cl_int})::cl_sampler end @checked function clRetainSampler(sampler) - @ccall libopencl.clRetainSampler(sampler::cl_sampler)::cl_int + @gcsafe_ccall libopencl.clRetainSampler(sampler::cl_sampler)::cl_int end @checked function clReleaseSampler(sampler) - @ccall libopencl.clReleaseSampler(sampler::cl_sampler)::cl_int + @gcsafe_ccall libopencl.clReleaseSampler(sampler::cl_sampler)::cl_int end @checked function clGetSamplerInfo(sampler, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetSamplerInfo(sampler::cl_sampler, param_name::cl_sampler_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetSamplerInfo(sampler::cl_sampler, + param_name::cl_sampler_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end function clCreateProgramWithSource(context, count, strings, lengths, errcode_ret) - @ccall libopencl.clCreateProgramWithSource(context::cl_context, count::cl_uint, - strings::Ptr{Ptr{Cchar}}, - lengths::Ptr{Csize_t}, - errcode_ret::Ptr{cl_int})::cl_program + @gcsafe_ccall libopencl.clCreateProgramWithSource(context::cl_context, count::cl_uint, + strings::Ptr{Ptr{Cchar}}, + lengths::Ptr{Csize_t}, + errcode_ret::Ptr{cl_int})::cl_program end function clCreateProgramWithBinary(context, num_devices, device_list, lengths, binaries, binary_status, errcode_ret) - @ccall libopencl.clCreateProgramWithBinary(context::cl_context, num_devices::cl_uint, - device_list::Ptr{cl_device_id}, - lengths::Ptr{Csize_t}, - binaries::Ptr{Ptr{Cuchar}}, - binary_status::Ptr{cl_int}, - errcode_ret::Ptr{cl_int})::cl_program + @gcsafe_ccall libopencl.clCreateProgramWithBinary(context::cl_context, + num_devices::cl_uint, + device_list::Ptr{cl_device_id}, + lengths::Ptr{Csize_t}, + binaries::Ptr{Ptr{Cuchar}}, + binary_status::Ptr{cl_int}, + errcode_ret::Ptr{cl_int})::cl_program end function clCreateProgramWithBuiltInKernels(context, num_devices, device_list, kernel_names, errcode_ret) - @ccall libopencl.clCreateProgramWithBuiltInKernels(context::cl_context, - num_devices::cl_uint, - device_list::Ptr{cl_device_id}, - kernel_names::Ptr{Cchar}, - errcode_ret::Ptr{cl_int})::cl_program + @gcsafe_ccall libopencl.clCreateProgramWithBuiltInKernels(context::cl_context, + num_devices::cl_uint, + device_list::Ptr{cl_device_id}, + kernel_names::Ptr{Cchar}, + errcode_ret::Ptr{cl_int})::cl_program end function clCreateProgramWithIL(context, il, length, errcode_ret) - @ccall libopencl.clCreateProgramWithIL(context::cl_context, il::Ptr{Cvoid}, - length::Csize_t, - errcode_ret::Ptr{cl_int})::cl_program + @gcsafe_ccall libopencl.clCreateProgramWithIL(context::cl_context, il::Ptr{Cvoid}, + length::Csize_t, + errcode_ret::Ptr{cl_int})::cl_program end @checked function clRetainProgram(program) - @ccall libopencl.clRetainProgram(program::cl_program)::cl_int + @gcsafe_ccall libopencl.clRetainProgram(program::cl_program)::cl_int end @checked function clReleaseProgram(program) - @ccall libopencl.clReleaseProgram(program::cl_program)::cl_int + @gcsafe_ccall libopencl.clReleaseProgram(program::cl_program)::cl_int end @checked function clBuildProgram(program, num_devices, device_list, options, pfn_notify, user_data) - @ccall libopencl.clBuildProgram(program::cl_program, num_devices::cl_uint, - device_list::Ptr{cl_device_id}, options::Ptr{Cchar}, - pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clBuildProgram(program::cl_program, num_devices::cl_uint, + device_list::Ptr{cl_device_id}, + options::Ptr{Cchar}, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid})::cl_int end @checked function clCompileProgram(program, num_devices, device_list, options, num_input_headers, input_headers, header_include_names, pfn_notify, user_data) - @ccall libopencl.clCompileProgram(program::cl_program, num_devices::cl_uint, - device_list::Ptr{cl_device_id}, options::Ptr{Cchar}, - num_input_headers::cl_uint, - input_headers::Ptr{cl_program}, - header_include_names::Ptr{Ptr{Cchar}}, - pfn_notify::Ptr{Cvoid}, user_data::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clCompileProgram(program::cl_program, num_devices::cl_uint, + device_list::Ptr{cl_device_id}, + options::Ptr{Cchar}, + num_input_headers::cl_uint, + input_headers::Ptr{cl_program}, + header_include_names::Ptr{Ptr{Cchar}}, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid})::cl_int end function clLinkProgram(context, num_devices, device_list, options, num_input_programs, input_programs, pfn_notify, user_data, errcode_ret) - @ccall libopencl.clLinkProgram(context::cl_context, num_devices::cl_uint, - device_list::Ptr{cl_device_id}, options::Ptr{Cchar}, - num_input_programs::cl_uint, - input_programs::Ptr{cl_program}, pfn_notify::Ptr{Cvoid}, - user_data::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_program + @gcsafe_ccall libopencl.clLinkProgram(context::cl_context, num_devices::cl_uint, + device_list::Ptr{cl_device_id}, + options::Ptr{Cchar}, + num_input_programs::cl_uint, + input_programs::Ptr{cl_program}, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_program end @checked function clSetProgramReleaseCallback(program, pfn_notify, user_data) - @ccall libopencl.clSetProgramReleaseCallback(program::cl_program, - pfn_notify::Ptr{Cvoid}, - user_data::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetProgramReleaseCallback(program::cl_program, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid})::cl_int end @checked function clSetProgramSpecializationConstant(program, spec_id, spec_size, spec_value) - @ccall libopencl.clSetProgramSpecializationConstant(program::cl_program, - spec_id::cl_uint, - spec_size::Csize_t, - spec_value::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetProgramSpecializationConstant(program::cl_program, + spec_id::cl_uint, + spec_size::Csize_t, + spec_value::Ptr{Cvoid})::cl_int end @checked function clUnloadPlatformCompiler(platform) - @ccall libopencl.clUnloadPlatformCompiler(platform::cl_platform_id)::cl_int + @gcsafe_ccall libopencl.clUnloadPlatformCompiler(platform::cl_platform_id)::cl_int end @checked function clGetProgramInfo(program, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetProgramInfo(program::cl_program, param_name::cl_program_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetProgramInfo(program::cl_program, + param_name::cl_program_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetProgramBuildInfo(program, device, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetProgramBuildInfo(program::cl_program, device::cl_device_id, - param_name::cl_program_build_info, - param_value_size::Csize_t, - param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetProgramBuildInfo(program::cl_program, device::cl_device_id, + param_name::cl_program_build_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end function clCreateKernel(program, kernel_name, errcode_ret) - @ccall libopencl.clCreateKernel(program::cl_program, kernel_name::Ptr{Cchar}, - errcode_ret::Ptr{cl_int})::cl_kernel + @gcsafe_ccall libopencl.clCreateKernel(program::cl_program, kernel_name::Ptr{Cchar}, + errcode_ret::Ptr{cl_int})::cl_kernel end @checked function clCreateKernelsInProgram(program, num_kernels, kernels, num_kernels_ret) - @ccall libopencl.clCreateKernelsInProgram(program::cl_program, num_kernels::cl_uint, - kernels::Ptr{cl_kernel}, - num_kernels_ret::Ptr{cl_uint})::cl_int + @gcsafe_ccall libopencl.clCreateKernelsInProgram(program::cl_program, + num_kernels::cl_uint, + kernels::Ptr{cl_kernel}, + num_kernels_ret::Ptr{cl_uint})::cl_int end function clCloneKernel(source_kernel, errcode_ret) - @ccall libopencl.clCloneKernel(source_kernel::cl_kernel, - errcode_ret::Ptr{cl_int})::cl_kernel + @gcsafe_ccall libopencl.clCloneKernel(source_kernel::cl_kernel, + errcode_ret::Ptr{cl_int})::cl_kernel end @checked function clRetainKernel(kernel) - @ccall libopencl.clRetainKernel(kernel::cl_kernel)::cl_int + @gcsafe_ccall libopencl.clRetainKernel(kernel::cl_kernel)::cl_int end @checked function clReleaseKernel(kernel) - @ccall libopencl.clReleaseKernel(kernel::cl_kernel)::cl_int + @gcsafe_ccall libopencl.clReleaseKernel(kernel::cl_kernel)::cl_int end @checked function clSetKernelArg(kernel, arg_index, arg_size, arg_value) - @ccall libopencl.clSetKernelArg(kernel::cl_kernel, arg_index::cl_uint, - arg_size::Csize_t, arg_value::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetKernelArg(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.clSetKernelArgSVMPointer(kernel::cl_kernel, arg_index::cl_uint, - arg_value::PtrOrCLPtr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetKernelArgSVMPointer(kernel::cl_kernel, arg_index::cl_uint, + arg_value::PtrOrCLPtr{Cvoid})::cl_int end @checked function clSetKernelExecInfo(kernel, param_name, param_value_size, param_value) - @ccall libopencl.clSetKernelExecInfo(kernel::cl_kernel, param_name::cl_kernel_exec_info, - param_value_size::Csize_t, - param_value::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetKernelExecInfo(kernel::cl_kernel, + param_name::cl_kernel_exec_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid})::cl_int end @checked function clGetKernelInfo(kernel, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetKernelInfo(kernel::cl_kernel, param_name::cl_kernel_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetKernelInfo(kernel::cl_kernel, param_name::cl_kernel_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetKernelArgInfo(kernel, arg_indx, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetKernelArgInfo(kernel::cl_kernel, arg_indx::cl_uint, - param_name::cl_kernel_arg_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetKernelArgInfo(kernel::cl_kernel, arg_indx::cl_uint, + param_name::cl_kernel_arg_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetKernelWorkGroupInfo(kernel, device, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetKernelWorkGroupInfo(kernel::cl_kernel, device::cl_device_id, - param_name::cl_kernel_work_group_info, - param_value_size::Csize_t, - param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetKernelWorkGroupInfo(kernel::cl_kernel, + device::cl_device_id, + param_name::cl_kernel_work_group_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clGetKernelSubGroupInfo(kernel, device, param_name, input_value_size, input_value, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetKernelSubGroupInfo(kernel::cl_kernel, device::cl_device_id, - param_name::cl_kernel_sub_group_info, - input_value_size::Csize_t, - input_value::Ptr{Cvoid}, - param_value_size::Csize_t, - param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetKernelSubGroupInfo(kernel::cl_kernel, device::cl_device_id, + param_name::cl_kernel_sub_group_info, + input_value_size::Csize_t, + input_value::Ptr{Cvoid}, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clWaitForEvents(num_events, event_list) - @ccall libopencl.clWaitForEvents(num_events::cl_uint, event_list::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clWaitForEvents(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.clGetEventInfo(event::cl_event, param_name::cl_event_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetEventInfo(event::cl_event, param_name::cl_event_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end function clCreateUserEvent(context, errcode_ret) - @ccall libopencl.clCreateUserEvent(context::cl_context, - errcode_ret::Ptr{cl_int})::cl_event + @gcsafe_ccall libopencl.clCreateUserEvent(context::cl_context, + errcode_ret::Ptr{cl_int})::cl_event end @checked function clRetainEvent(event) - @ccall libopencl.clRetainEvent(event::cl_event)::cl_int + @gcsafe_ccall libopencl.clRetainEvent(event::cl_event)::cl_int end @checked function clReleaseEvent(event) - @ccall libopencl.clReleaseEvent(event::cl_event)::cl_int + @gcsafe_ccall libopencl.clReleaseEvent(event::cl_event)::cl_int end @checked function clSetUserEventStatus(event, execution_status) - @ccall libopencl.clSetUserEventStatus(event::cl_event, execution_status::cl_int)::cl_int + @gcsafe_ccall libopencl.clSetUserEventStatus(event::cl_event, + execution_status::cl_int)::cl_int end @checked function clSetEventCallback(event, command_exec_callback_type, pfn_notify, user_data) - @ccall libopencl.clSetEventCallback(event::cl_event, command_exec_callback_type::cl_int, - pfn_notify::Ptr{Cvoid}, - user_data::Ptr{Cvoid})::cl_int + @gcsafe_ccall libopencl.clSetEventCallback(event::cl_event, + command_exec_callback_type::cl_int, + pfn_notify::Ptr{Cvoid}, + user_data::Ptr{Cvoid})::cl_int end @checked function clGetEventProfilingInfo(event, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetEventProfilingInfo(event::cl_event, param_name::cl_profiling_info, - param_value_size::Csize_t, - param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetEventProfilingInfo(event::cl_event, + param_name::cl_profiling_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clFlush(command_queue) - @ccall libopencl.clFlush(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.clFlush(command_queue::cl_command_queue)::cl_int end @checked function clFinish(command_queue) - @ccall libopencl.clFinish(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.clFinish(command_queue::cl_command_queue)::cl_int end @checked function clEnqueueReadBuffer(command_queue, buffer, blocking_read, offset, size, ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueReadBuffer(command_queue::cl_command_queue, buffer::cl_mem, - blocking_read::cl_bool, offset::Csize_t, - size::Csize_t, ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueReadBuffer(command_queue::cl_command_queue, + buffer::cl_mem, + blocking_read::cl_bool, offset::Csize_t, + size::Csize_t, ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueReadBufferRect(command_queue, buffer, blocking_read, @@ -718,28 +759,30 @@ end buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueReadBufferRect(command_queue::cl_command_queue, - buffer::cl_mem, blocking_read::cl_bool, - buffer_origin::Ptr{Csize_t}, - host_origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, - buffer_row_pitch::Csize_t, - buffer_slice_pitch::Csize_t, - host_row_pitch::Csize_t, - host_slice_pitch::Csize_t, ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueReadBufferRect(command_queue::cl_command_queue, + buffer::cl_mem, blocking_read::cl_bool, + buffer_origin::Ptr{Csize_t}, + host_origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + buffer_row_pitch::Csize_t, + buffer_slice_pitch::Csize_t, + host_row_pitch::Csize_t, + host_slice_pitch::Csize_t, + ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueWriteBuffer(command_queue, buffer, blocking_write, offset, size, ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueWriteBuffer(command_queue::cl_command_queue, buffer::cl_mem, - blocking_write::cl_bool, offset::Csize_t, - size::Csize_t, ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueWriteBuffer(command_queue::cl_command_queue, + buffer::cl_mem, + blocking_write::cl_bool, offset::Csize_t, + size::Csize_t, ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueWriteBufferRect(command_queue, buffer, blocking_write, @@ -747,349 +790,375 @@ end buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueWriteBufferRect(command_queue::cl_command_queue, - buffer::cl_mem, blocking_write::cl_bool, - buffer_origin::Ptr{Csize_t}, - host_origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, - buffer_row_pitch::Csize_t, - buffer_slice_pitch::Csize_t, - host_row_pitch::Csize_t, - host_slice_pitch::Csize_t, ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueWriteBufferRect(command_queue::cl_command_queue, + buffer::cl_mem, + blocking_write::cl_bool, + buffer_origin::Ptr{Csize_t}, + host_origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + buffer_row_pitch::Csize_t, + buffer_slice_pitch::Csize_t, + host_row_pitch::Csize_t, + host_slice_pitch::Csize_t, + ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueFillBuffer(command_queue, buffer, pattern, pattern_size, offset, size, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueFillBuffer(command_queue::cl_command_queue, buffer::cl_mem, - pattern::Ptr{Cvoid}, pattern_size::Csize_t, - offset::Csize_t, size::Csize_t, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueFillBuffer(command_queue::cl_command_queue, + buffer::cl_mem, + pattern::Ptr{Cvoid}, pattern_size::Csize_t, + offset::Csize_t, size::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueCopyBuffer(command_queue, src_buffer, dst_buffer, src_offset, dst_offset, size, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueCopyBuffer(command_queue::cl_command_queue, - src_buffer::cl_mem, dst_buffer::cl_mem, - src_offset::Csize_t, dst_offset::Csize_t, - size::Csize_t, num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueCopyBuffer(command_queue::cl_command_queue, + src_buffer::cl_mem, dst_buffer::cl_mem, + src_offset::Csize_t, dst_offset::Csize_t, + size::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueCopyBufferRect(command_queue, src_buffer, dst_buffer, src_origin, dst_origin, region, src_row_pitch, src_slice_pitch, dst_row_pitch, dst_slice_pitch, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueCopyBufferRect(command_queue::cl_command_queue, - src_buffer::cl_mem, dst_buffer::cl_mem, - src_origin::Ptr{Csize_t}, - dst_origin::Ptr{Csize_t}, region::Ptr{Csize_t}, - src_row_pitch::Csize_t, - src_slice_pitch::Csize_t, - dst_row_pitch::Csize_t, - dst_slice_pitch::Csize_t, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueCopyBufferRect(command_queue::cl_command_queue, + src_buffer::cl_mem, dst_buffer::cl_mem, + src_origin::Ptr{Csize_t}, + dst_origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + src_row_pitch::Csize_t, + src_slice_pitch::Csize_t, + dst_row_pitch::Csize_t, + dst_slice_pitch::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueReadImage(command_queue, image, blocking_read, origin, region, row_pitch, slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueReadImage(command_queue::cl_command_queue, image::cl_mem, - blocking_read::cl_bool, origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, row_pitch::Csize_t, - slice_pitch::Csize_t, ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueReadImage(command_queue::cl_command_queue, + image::cl_mem, + blocking_read::cl_bool, origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, row_pitch::Csize_t, + slice_pitch::Csize_t, ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueWriteImage(command_queue, image, blocking_write, origin, region, input_row_pitch, input_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueWriteImage(command_queue::cl_command_queue, image::cl_mem, - blocking_write::cl_bool, origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, input_row_pitch::Csize_t, - input_slice_pitch::Csize_t, ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueWriteImage(command_queue::cl_command_queue, + image::cl_mem, + blocking_write::cl_bool, + origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + input_row_pitch::Csize_t, + input_slice_pitch::Csize_t, ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueFillImage(command_queue, image, fill_color, origin, region, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueFillImage(command_queue::cl_command_queue, image::cl_mem, - fill_color::Ptr{Cvoid}, origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueFillImage(command_queue::cl_command_queue, + image::cl_mem, + fill_color::Ptr{Cvoid}, origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueCopyImage(command_queue, src_image, dst_image, src_origin, dst_origin, region, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueCopyImage(command_queue::cl_command_queue, src_image::cl_mem, - dst_image::cl_mem, src_origin::Ptr{Csize_t}, - dst_origin::Ptr{Csize_t}, region::Ptr{Csize_t}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueCopyImage(command_queue::cl_command_queue, + src_image::cl_mem, + dst_image::cl_mem, src_origin::Ptr{Csize_t}, + dst_origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueCopyImageToBuffer(command_queue, src_image, dst_buffer, src_origin, region, dst_offset, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueCopyImageToBuffer(command_queue::cl_command_queue, - src_image::cl_mem, dst_buffer::cl_mem, - src_origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, dst_offset::Csize_t, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueCopyImageToBuffer(command_queue::cl_command_queue, + src_image::cl_mem, + dst_buffer::cl_mem, + src_origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + dst_offset::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueCopyBufferToImage(command_queue, src_buffer, dst_image, src_offset, dst_origin, region, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueCopyBufferToImage(command_queue::cl_command_queue, - src_buffer::cl_mem, dst_image::cl_mem, - src_offset::Csize_t, - dst_origin::Ptr{Csize_t}, - region::Ptr{Csize_t}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueCopyBufferToImage(command_queue::cl_command_queue, + src_buffer::cl_mem, + dst_image::cl_mem, + src_offset::Csize_t, + dst_origin::Ptr{Csize_t}, + region::Ptr{Csize_t}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end function clEnqueueMapBuffer(command_queue, buffer, blocking_map, map_flags, offset, size, num_events_in_wait_list, event_wait_list, event, errcode_ret) - @ccall libopencl.clEnqueueMapBuffer(command_queue::cl_command_queue, buffer::cl_mem, - blocking_map::cl_bool, map_flags::cl_map_flags, - offset::Csize_t, size::Csize_t, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event}, - errcode_ret::Ptr{cl_int})::Ptr{Cvoid} + @gcsafe_ccall libopencl.clEnqueueMapBuffer(command_queue::cl_command_queue, + buffer::cl_mem, + blocking_map::cl_bool, + map_flags::cl_map_flags, + offset::Csize_t, size::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event}, + errcode_ret::Ptr{cl_int})::Ptr{Cvoid} end function clEnqueueMapImage(command_queue, image, blocking_map, map_flags, origin, region, image_row_pitch, image_slice_pitch, num_events_in_wait_list, event_wait_list, event, errcode_ret) - @ccall libopencl.clEnqueueMapImage(command_queue::cl_command_queue, image::cl_mem, - blocking_map::cl_bool, map_flags::cl_map_flags, - origin::Ptr{Csize_t}, region::Ptr{Csize_t}, - image_row_pitch::Ptr{Csize_t}, - image_slice_pitch::Ptr{Csize_t}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, event::Ptr{cl_event}, - errcode_ret::Ptr{cl_int})::Ptr{Cvoid} + @gcsafe_ccall libopencl.clEnqueueMapImage(command_queue::cl_command_queue, + image::cl_mem, + blocking_map::cl_bool, + map_flags::cl_map_flags, + origin::Ptr{Csize_t}, region::Ptr{Csize_t}, + image_row_pitch::Ptr{Csize_t}, + image_slice_pitch::Ptr{Csize_t}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event}, + errcode_ret::Ptr{cl_int})::Ptr{Cvoid} end @checked function clEnqueueUnmapMemObject(command_queue, memobj, mapped_ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueUnmapMemObject(command_queue::cl_command_queue, - memobj::cl_mem, mapped_ptr::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueUnmapMemObject(command_queue::cl_command_queue, + memobj::cl_mem, mapped_ptr::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueMigrateMemObjects(command_queue, num_mem_objects, mem_objects, flags, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueMigrateMemObjects(command_queue::cl_command_queue, - num_mem_objects::cl_uint, - mem_objects::Ptr{cl_mem}, - flags::cl_mem_migration_flags, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueMigrateMemObjects(command_queue::cl_command_queue, + num_mem_objects::cl_uint, + mem_objects::Ptr{cl_mem}, + flags::cl_mem_migration_flags, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueNDRangeKernel(command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueNDRangeKernel(command_queue::cl_command_queue, - kernel::cl_kernel, work_dim::cl_uint, - global_work_offset::Ptr{Csize_t}, - global_work_size::Ptr{Csize_t}, - local_work_size::Ptr{Csize_t}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueNDRangeKernel(command_queue::cl_command_queue, + kernel::cl_kernel, work_dim::cl_uint, + global_work_offset::Ptr{Csize_t}, + global_work_size::Ptr{Csize_t}, + local_work_size::Ptr{Csize_t}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueNativeKernel(command_queue, user_func, args, cb_args, num_mem_objects, mem_list, args_mem_loc, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueNativeKernel(command_queue::cl_command_queue, - user_func::Ptr{Cvoid}, args::Ptr{Cvoid}, - cb_args::Csize_t, num_mem_objects::cl_uint, - mem_list::Ptr{cl_mem}, - args_mem_loc::Ptr{Ptr{Cvoid}}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueNativeKernel(command_queue::cl_command_queue, + user_func::Ptr{Cvoid}, args::Ptr{Cvoid}, + cb_args::Csize_t, + num_mem_objects::cl_uint, + mem_list::Ptr{cl_mem}, + args_mem_loc::Ptr{Ptr{Cvoid}}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueMarkerWithWaitList(command_queue, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueMarkerWithWaitList(command_queue::cl_command_queue, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueMarkerWithWaitList(command_queue::cl_command_queue, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueBarrierWithWaitList(command_queue, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueBarrierWithWaitList(command_queue::cl_command_queue, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueBarrierWithWaitList(command_queue::cl_command_queue, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueSVMFree(command_queue, num_svm_pointers, svm_pointers, pfn_free_func, user_data, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueSVMFree(command_queue::cl_command_queue, - num_svm_pointers::cl_uint, - svm_pointers::Ptr{Ptr{Cvoid}}, - pfn_free_func::Ptr{Cvoid}, user_data::Ptr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueSVMFree(command_queue::cl_command_queue, + num_svm_pointers::cl_uint, + svm_pointers::Ptr{Ptr{Cvoid}}, + pfn_free_func::Ptr{Cvoid}, + user_data::Ptr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueSVMMemcpy(command_queue, blocking_copy, dst_ptr, src_ptr, size, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueSVMMemcpy(command_queue::cl_command_queue, - blocking_copy::cl_bool, dst_ptr::PtrOrCLPtr{Cvoid}, - src_ptr::PtrOrCLPtr{Cvoid}, size::Csize_t, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueSVMMemcpy(command_queue::cl_command_queue, + blocking_copy::cl_bool, + dst_ptr::PtrOrCLPtr{Cvoid}, + src_ptr::PtrOrCLPtr{Cvoid}, size::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueSVMMemFill(command_queue, svm_ptr, pattern, pattern_size, size, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueSVMMemFill(command_queue::cl_command_queue, - svm_ptr::CLPtr{Cvoid}, pattern::Ptr{Cvoid}, - pattern_size::Csize_t, size::Csize_t, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueSVMMemFill(command_queue::cl_command_queue, + svm_ptr::CLPtr{Cvoid}, pattern::Ptr{Cvoid}, + pattern_size::Csize_t, size::Csize_t, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueSVMMap(command_queue, blocking_map, flags, svm_ptr, size, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueSVMMap(command_queue::cl_command_queue, blocking_map::cl_bool, - flags::cl_map_flags, svm_ptr::CLPtr{Cvoid}, - size::Csize_t, num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueSVMMap(command_queue::cl_command_queue, + blocking_map::cl_bool, + flags::cl_map_flags, svm_ptr::CLPtr{Cvoid}, + size::Csize_t, num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueSVMUnmap(command_queue, svm_ptr, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueSVMUnmap(command_queue::cl_command_queue, - svm_ptr::CLPtr{Cvoid}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueSVMUnmap(command_queue::cl_command_queue, + svm_ptr::CLPtr{Cvoid}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueSVMMigrateMem(command_queue, num_svm_pointers, svm_pointers, sizes, flags, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueSVMMigrateMem(command_queue::cl_command_queue, - num_svm_pointers::cl_uint, - svm_pointers::Ptr{CLPtr{Cvoid}}, - sizes::Ptr{Csize_t}, - flags::cl_mem_migration_flags, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueSVMMigrateMem(command_queue::cl_command_queue, + num_svm_pointers::cl_uint, + svm_pointers::Ptr{CLPtr{Cvoid}}, + sizes::Ptr{Csize_t}, + flags::cl_mem_migration_flags, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end function clGetExtensionFunctionAddressForPlatform(platform, func_name) - @ccall libopencl.clGetExtensionFunctionAddressForPlatform(platform::cl_platform_id, - func_name::Ptr{Cchar})::Ptr{Cvoid} + @gcsafe_ccall libopencl.clGetExtensionFunctionAddressForPlatform(platform::cl_platform_id, + func_name::Ptr{Cchar})::Ptr{Cvoid} end function clCreateImage2D(context, flags, image_format, image_width, image_height, image_row_pitch, host_ptr, errcode_ret) - @ccall libopencl.clCreateImage2D(context::cl_context, flags::cl_mem_flags, - image_format::Ptr{cl_image_format}, - image_width::Csize_t, image_height::Csize_t, - image_row_pitch::Csize_t, host_ptr::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateImage2D(context::cl_context, flags::cl_mem_flags, + image_format::Ptr{cl_image_format}, + image_width::Csize_t, image_height::Csize_t, + image_row_pitch::Csize_t, host_ptr::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreateImage3D(context, flags, image_format, image_width, image_height, image_depth, image_row_pitch, image_slice_pitch, host_ptr, errcode_ret) - @ccall libopencl.clCreateImage3D(context::cl_context, flags::cl_mem_flags, - image_format::Ptr{cl_image_format}, - image_width::Csize_t, image_height::Csize_t, - image_depth::Csize_t, image_row_pitch::Csize_t, - image_slice_pitch::Csize_t, host_ptr::Ptr{Cvoid}, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateImage3D(context::cl_context, flags::cl_mem_flags, + image_format::Ptr{cl_image_format}, + image_width::Csize_t, image_height::Csize_t, + image_depth::Csize_t, image_row_pitch::Csize_t, + image_slice_pitch::Csize_t, + host_ptr::Ptr{Cvoid}, + errcode_ret::Ptr{cl_int})::cl_mem end @checked function clEnqueueMarker(command_queue, event) - @ccall libopencl.clEnqueueMarker(command_queue::cl_command_queue, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueMarker(command_queue::cl_command_queue, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueWaitForEvents(command_queue, num_events, event_list) - @ccall libopencl.clEnqueueWaitForEvents(command_queue::cl_command_queue, - num_events::cl_uint, - event_list::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueWaitForEvents(command_queue::cl_command_queue, + num_events::cl_uint, + event_list::Ptr{cl_event})::cl_int end @checked function clEnqueueBarrier(command_queue) - @ccall libopencl.clEnqueueBarrier(command_queue::cl_command_queue)::cl_int + @gcsafe_ccall libopencl.clEnqueueBarrier(command_queue::cl_command_queue)::cl_int end @checked function clUnloadCompiler() - @ccall libopencl.clUnloadCompiler()::cl_int + @gcsafe_ccall libopencl.clUnloadCompiler()::cl_int end function clGetExtensionFunctionAddress(func_name) - @ccall libopencl.clGetExtensionFunctionAddress(func_name::Ptr{Cchar})::Ptr{Cvoid} + @gcsafe_ccall libopencl.clGetExtensionFunctionAddress(func_name::Ptr{Cchar})::Ptr{Cvoid} end function clCreateCommandQueue(context, device, properties, errcode_ret) - @ccall libopencl.clCreateCommandQueue(context::cl_context, device::cl_device_id, - properties::cl_command_queue_properties, - errcode_ret::Ptr{cl_int})::cl_command_queue + @gcsafe_ccall libopencl.clCreateCommandQueue(context::cl_context, device::cl_device_id, + properties::cl_command_queue_properties, + errcode_ret::Ptr{cl_int})::cl_command_queue end function clCreateSampler(context, normalized_coords, addressing_mode, filter_mode, errcode_ret) - @ccall libopencl.clCreateSampler(context::cl_context, normalized_coords::cl_bool, - addressing_mode::cl_addressing_mode, - filter_mode::cl_filter_mode, - errcode_ret::Ptr{cl_int})::cl_sampler + @gcsafe_ccall libopencl.clCreateSampler(context::cl_context, normalized_coords::cl_bool, + addressing_mode::cl_addressing_mode, + filter_mode::cl_filter_mode, + errcode_ret::Ptr{cl_int})::cl_sampler end @checked function clEnqueueTask(command_queue, kernel, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueTask(command_queue::cl_command_queue, kernel::cl_kernel, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueTask(command_queue::cl_command_queue, + kernel::cl_kernel, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end const cl_GLint = Cint @@ -1118,17 +1187,17 @@ const clCreateFromGLBuffer_fn = Ptr{clCreateFromGLBuffer_t} @checked function clGetGLContextInfoKHR(properties, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetGLContextInfoKHR(properties::Ptr{cl_context_properties}, - param_name::cl_gl_context_info, - param_value_size::Csize_t, - param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetGLContextInfoKHR(properties::Ptr{cl_context_properties}, + param_name::cl_gl_context_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end function clCreateFromGLBuffer(context, flags, bufobj, errcode_ret) - @ccall libopencl.clCreateFromGLBuffer(context::cl_context, flags::cl_mem_flags, - bufobj::cl_GLuint, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateFromGLBuffer(context::cl_context, flags::cl_mem_flags, + bufobj::cl_GLuint, + errcode_ret::Ptr{cl_int})::cl_mem end # typedef cl_mem CL_API_CALL clCreateFromGLTexture_t ( cl_context context , cl_mem_flags flags , cl_GLenum target , cl_GLint miplevel , cl_GLuint texture , cl_int * errcode_ret ) @@ -1137,10 +1206,10 @@ const clCreateFromGLTexture_t = Cvoid const clCreateFromGLTexture_fn = Ptr{clCreateFromGLTexture_t} function clCreateFromGLTexture(context, flags, target, miplevel, texture, errcode_ret) - @ccall libopencl.clCreateFromGLTexture(context::cl_context, flags::cl_mem_flags, - target::cl_GLenum, miplevel::cl_GLint, - texture::cl_GLuint, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateFromGLTexture(context::cl_context, flags::cl_mem_flags, + target::cl_GLenum, miplevel::cl_GLint, + texture::cl_GLuint, + errcode_ret::Ptr{cl_int})::cl_mem end # typedef cl_mem CL_API_CALL clCreateFromGLRenderbuffer_t ( cl_context context , cl_mem_flags flags , cl_GLuint renderbuffer , cl_int * errcode_ret ) @@ -1169,42 +1238,45 @@ const clEnqueueReleaseGLObjects_t = Cvoid const clEnqueueReleaseGLObjects_fn = Ptr{clEnqueueReleaseGLObjects_t} function clCreateFromGLRenderbuffer(context, flags, renderbuffer, errcode_ret) - @ccall libopencl.clCreateFromGLRenderbuffer(context::cl_context, flags::cl_mem_flags, - renderbuffer::cl_GLuint, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateFromGLRenderbuffer(context::cl_context, + flags::cl_mem_flags, + renderbuffer::cl_GLuint, + errcode_ret::Ptr{cl_int})::cl_mem end @checked function clGetGLObjectInfo(memobj, gl_object_type, gl_object_name) - @ccall libopencl.clGetGLObjectInfo(memobj::cl_mem, - gl_object_type::Ptr{cl_gl_object_type}, - gl_object_name::Ptr{cl_GLuint})::cl_int + @gcsafe_ccall libopencl.clGetGLObjectInfo(memobj::cl_mem, + gl_object_type::Ptr{cl_gl_object_type}, + gl_object_name::Ptr{cl_GLuint})::cl_int end @checked function clGetGLTextureInfo(memobj, param_name, param_value_size, param_value, param_value_size_ret) - @ccall libopencl.clGetGLTextureInfo(memobj::cl_mem, param_name::cl_gl_texture_info, - param_value_size::Csize_t, param_value::Ptr{Cvoid}, - param_value_size_ret::Ptr{Csize_t})::cl_int + @gcsafe_ccall libopencl.clGetGLTextureInfo(memobj::cl_mem, + param_name::cl_gl_texture_info, + param_value_size::Csize_t, + param_value::Ptr{Cvoid}, + param_value_size_ret::Ptr{Csize_t})::cl_int end @checked function clEnqueueAcquireGLObjects(command_queue, num_objects, mem_objects, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueAcquireGLObjects(command_queue::cl_command_queue, - num_objects::cl_uint, - mem_objects::Ptr{cl_mem}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueAcquireGLObjects(command_queue::cl_command_queue, + num_objects::cl_uint, + mem_objects::Ptr{cl_mem}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end @checked function clEnqueueReleaseGLObjects(command_queue, num_objects, mem_objects, num_events_in_wait_list, event_wait_list, event) - @ccall libopencl.clEnqueueReleaseGLObjects(command_queue::cl_command_queue, - num_objects::cl_uint, - mem_objects::Ptr{cl_mem}, - num_events_in_wait_list::cl_uint, - event_wait_list::Ptr{cl_event}, - event::Ptr{cl_event})::cl_int + @gcsafe_ccall libopencl.clEnqueueReleaseGLObjects(command_queue::cl_command_queue, + num_objects::cl_uint, + mem_objects::Ptr{cl_mem}, + num_events_in_wait_list::cl_uint, + event_wait_list::Ptr{cl_event}, + event::Ptr{cl_event})::cl_int end # typedef cl_mem CL_API_CALL clCreateFromGLTexture2D_t ( cl_context context , cl_mem_flags flags , cl_GLenum target , cl_GLint miplevel , cl_GLuint texture , cl_int * errcode_ret ) @@ -1218,17 +1290,19 @@ const clCreateFromGLTexture3D_t = Cvoid const clCreateFromGLTexture3D_fn = Ptr{clCreateFromGLTexture3D_t} function clCreateFromGLTexture2D(context, flags, target, miplevel, texture, errcode_ret) - @ccall libopencl.clCreateFromGLTexture2D(context::cl_context, flags::cl_mem_flags, - target::cl_GLenum, miplevel::cl_GLint, - texture::cl_GLuint, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateFromGLTexture2D(context::cl_context, + flags::cl_mem_flags, + target::cl_GLenum, miplevel::cl_GLint, + texture::cl_GLuint, + errcode_ret::Ptr{cl_int})::cl_mem end function clCreateFromGLTexture3D(context, flags, target, miplevel, texture, errcode_ret) - @ccall libopencl.clCreateFromGLTexture3D(context::cl_context, flags::cl_mem_flags, - target::cl_GLenum, miplevel::cl_GLint, - texture::cl_GLuint, - errcode_ret::Ptr{cl_int})::cl_mem + @gcsafe_ccall libopencl.clCreateFromGLTexture3D(context::cl_context, + flags::cl_mem_flags, + target::cl_GLenum, miplevel::cl_GLint, + texture::cl_GLuint, + errcode_ret::Ptr{cl_int})::cl_mem end mutable struct __GLsync end @@ -1241,8 +1315,8 @@ const clCreateEventFromGLsyncKHR_t = Cvoid const clCreateEventFromGLsyncKHR_fn = Ptr{clCreateEventFromGLsyncKHR_t} function clCreateEventFromGLsyncKHR(context, sync, errcode_ret) - @ccall libopencl.clCreateEventFromGLsyncKHR(context::cl_context, sync::cl_GLsync, - errcode_ret::Ptr{cl_int})::cl_event + @gcsafe_ccall libopencl.clCreateEventFromGLsyncKHR(context::cl_context, sync::cl_GLsync, + errcode_ret::Ptr{cl_int})::cl_event end # typedef cl_int CL_API_CALL clGetSupportedGLTextureFormatsINTEL_t ( cl_context context , cl_mem_flags flags , cl_mem_object_type image_type , cl_uint num_entries , cl_GLenum * gl_formats , cl_uint * num_texture_formats ) @@ -1253,12 +1327,12 @@ const clGetSupportedGLTextureFormatsINTEL_fn = Ptr{clGetSupportedGLTextureFormat @checked function clGetSupportedGLTextureFormatsINTEL(context, flags, image_type, num_entries, gl_formats, num_texture_formats) - @ccall libopencl.clGetSupportedGLTextureFormatsINTEL(context::cl_context, - flags::cl_mem_flags, - image_type::cl_mem_object_type, - num_entries::cl_uint, - gl_formats::Ptr{cl_GLenum}, - num_texture_formats::Ptr{cl_uint})::cl_int + @gcsafe_ccall libopencl.clGetSupportedGLTextureFormatsINTEL(context::cl_context, + flags::cl_mem_flags, + image_type::cl_mem_object_type, + num_entries::cl_uint, + gl_formats::Ptr{cl_GLenum}, + num_texture_formats::Ptr{cl_uint})::cl_int end const cl_device_partition_property_ext = cl_ulong diff --git a/lib/cl/pointer.jl b/lib/cl/pointer.jl index a78146b0..179b32aa 100644 --- a/lib/cl/pointer.jl +++ b/lib/cl/pointer.jl @@ -138,6 +138,10 @@ function Base.unsafe_convert(::Type{PtrOrCLPtr{T}}, val) where {T} return Base.bitcast(PtrOrCLPtr{T}, ptr) end +# `@gcsafe_ccall` converts arguments itself before passing them to `ccall`, which converts +# them again, so make that second (identity) conversion unambiguous +Base.unsafe_convert(::Type{PtrOrCLPtr{T}}, x::PtrOrCLPtr{T}) where {T} = x + # # Device reference objects @@ -207,6 +211,8 @@ Base.unsafe_convert(::Type{RefOrCLRef{T}}, x::Ref{T}) where {T} = Base.bitcast(RefOrCLRef{T}, Base.unsafe_convert(Ptr{T}, x)) Base.unsafe_convert(::Type{RefOrCLRef{T}}, x) where {T} = Base.bitcast(RefOrCLRef{T}, Base.unsafe_convert(Ptr{T}, x)) +# see `PtrOrCLPtr` +Base.unsafe_convert(::Type{RefOrCLRef{T}}, x::RefOrCLRef{T}) where {T} = x # support conversion from GPU ref Base.unsafe_convert(::Type{RefOrCLRef{T}}, x::CLRefs{T}) where {T} = diff --git a/res/wrap.jl b/res/wrap.jl index 030cba73..5c479b97 100644 --- a/res/wrap.jl +++ b/res/wrap.jl @@ -77,9 +77,14 @@ function rewriter!(ctx, options) target_expr = call_expr.args[1].args[1] fn = String(target_expr.args[2].value) - # replace `@ccall` with `@ext_ccall` when dealing with an extension function + # replace `@ccall` with `@ext_ccall` when dealing with an extension function, + # and with `@gcsafe_ccall` otherwise: API calls may block (waiting for the + # device, compiling, ...), which should not prevent other threads from + # collecting garbage. if basename(get_filename(node.cursor)) == "cl_ext.h" expr.args[2].args[1].args[1] = Symbol("@ext_ccall") + else + expr.args[2].args[1].args[1] = Symbol("@gcsafe_ccall") end # rewrite pointer return types From c3def13ab5d1163ebe5d33278883769190659334 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 21:23:32 +0200 Subject: [PATCH 2/7] Wait for events cooperatively `wait(::cl.AbstractEvent)` called `clWaitForEvents`, blocking the calling thread until the commands completed. No other task could run on that thread in the meantime, and waiting could not be interrupted. Use GPUToolbox's `cooperative_wait` instead, like CUDA.jl does. It polls the execution status of the events for a while, which keeps short waits fast, and then hands a blocking `clWaitForEvents` to a worker thread while the task yields. Since polling only makes progress once commands have been submitted, the events' queues are flushed first. Afterwards, `clWaitForEvents` is called again (returning immediately) to synchronize host memory and to report errors. Waiting falls back to blocking in finalizers and while precompiling, and the `nonblocking_synchronization` preference switches back to blocking entirely. --- LocalPreferences.toml | 5 +++ lib/cl/CL.jl | 1 + lib/cl/api.jl | 2 +- lib/cl/event.jl | 20 ++-------- lib/cl/synchronization.jl | 80 +++++++++++++++++++++++++++++++++++++++ src/util.jl | 1 + test/synchronization.jl | 62 ++++++++++++++++++++++++++++++ 7 files changed, 153 insertions(+), 18 deletions(-) create mode 100644 lib/cl/synchronization.jl diff --git a/LocalPreferences.toml b/LocalPreferences.toml index 3177d27e..53ed1823 100644 --- a/LocalPreferences.toml +++ b/LocalPreferences.toml @@ -11,3 +11,8 @@ # - "khronos": the Khronos SPIRV-LLVM-Translator, which additionally requires # `SPIRV_LLVM_Translator_jll` (or `SPIRV_LLVM_Translator_unified_jll`) to be loaded #llvm_to_spirv_backend="llvm" + +# Whether to wait for the device cooperatively, yielding to other tasks instead of blocking +# the calling thread in the OpenCL driver. Disable for bisecting issues, or to compare +# against the blocking behavior. +#nonblocking_synchronization=true diff --git a/lib/cl/CL.jl b/lib/cl/CL.jl index 2788a72f..6000f8de 100644 --- a/lib/cl/CL.jl +++ b/lib/cl/CL.jl @@ -20,6 +20,7 @@ include("device.jl") include("context.jl") include("cmdqueue.jl") include("event.jl") +include("synchronization.jl") include("memory.jl") include("program.jl") include("kernel.jl") diff --git a/lib/cl/api.jl b/lib/cl/api.jl index cd2d3288..86721ca9 100644 --- a/lib/cl/api.jl +++ b/lib/cl/api.jl @@ -1,5 +1,5 @@ import OpenCL_jll -using GPUToolbox: @gcsafe_ccall +using GPUToolbox: @gcsafe_ccall, cooperative_wait const libopencl = OpenCL_jll.libopencl diff --git a/lib/cl/event.jl b/lib/cl/event.jl index 67a53010..e0462959 100644 --- a/lib/cl/event.jl +++ b/lib/cl/event.jl @@ -143,26 +143,12 @@ function add_callback(evt::AbstractEvent, callback::Function) end end -function Base.wait(evt::AbstractEvent) - evt_id = [evt.id] - clWaitForEvents(cl_uint(1), evt_id) - return evt -end - -function Base.wait(evts::Vector{AbstractEvent}) - isempty(evts) && return evts - evt_ids = [pointer(evt) for evt in evts] - GC.@preserve evts begin - clWaitForEvents(length(evt_ids), evt_ids) - end - return evts -end - -function enqueue_marker_with_wait_list(wait_for::Vector{AbstractEvent}) +function enqueue_marker_with_wait_list(wait_for::Vector{AbstractEvent}; + queue::CmdQueue=queue()) n_wait_events = cl_uint(length(wait_for)) wait_evt_ids = [evt.id for evt in wait_for] ret_evt = Ref{cl_event}() - clEnqueueMarkerWithWaitList(queue(), n_wait_events, + clEnqueueMarkerWithWaitList(queue, n_wait_events, isempty(wait_evt_ids) ? C_NULL : wait_evt_ids, ret_evt) @return_event ret_evt[] diff --git a/lib/cl/synchronization.jl b/lib/cl/synchronization.jl new file mode 100644 index 00000000..2ea032cc --- /dev/null +++ b/lib/cl/synchronization.jl @@ -0,0 +1,80 @@ +# cooperative synchronization +# +# Waiting for the device should not block the calling thread in the OpenCL driver, but +# yield to the Julia scheduler so that other tasks can run in the meantime. GPUToolbox's +# `cooperative_wait` implements this by polling the status of commands for a while, and +# then handing the blocking wait to a worker thread. + +# whether to wait for the device cooperatively. disable to block in the OpenCL driver +# instead, e.g., for bisecting issues or comparing against the blocking behavior. +const nonblocking_synchronization = + @load_preference("nonblocking_synchronization", true) + +event_ids(evts) = cl_event[pointer(evt) for evt in evts] + +# commands have completed when their execution status is `CL_COMPLETE` (zero), or negative +# when they were terminated abnormally +function iscomplete(evt::AbstractEvent) + status = Ref{Cint}() + clGetEventInfo(evt, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Cint), status, C_NULL) + return status[] <= CL_COMPLETE +end + +# commands only need to start executing once the queue they were submitted to has been +# flushed. waiting in the driver flushes implicitly, but polling does not. +function flush_queue(evt::AbstractEvent) + queue = Ref{cl_command_queue}() + clGetEventInfo(evt, CL_EVENT_COMMAND_QUEUE, sizeof(cl_command_queue), queue, C_NULL) + # user events do not belong to a queue + queue[] == C_NULL || clFlush(queue[]) + return +end + +# wait for events to complete, throwing a `CLError` if a command was terminated abnormally. +# +# by default, the wait can be interrupted, in which case the commands may still be +# executing. waits that cannot be interrupted only return (or throw) once the commands have +# completed, which is required when they access memory that the caller releases upon +# returning. +function wait_events(evts::Vector{<:AbstractEvent}; cancellable::Bool=true) + if nonblocking_synchronization + try + foreach(flush_queue, evts) + cooperative_wait(evts; isdone=evts -> all(iscomplete, evts), cancellable) do evts + GC.@preserve evts unchecked_clWaitForEvents(length(evts), event_ids(evts)) + end + catch + # waits that cannot be interrupted only throw once the commands have completed + # (or waiting failed, in which case this blocks), but host memory still needs to + # be synchronized before the caller unwinds + cancellable || GC.@preserve evts begin + unchecked_clWaitForEvents(length(evts), event_ids(evts)) + end + rethrow() + end + end + + # synchronize host memory and report errors (without blocking, if we waited above) + GC.@preserve evts clWaitForEvents(length(evts), event_ids(evts)) + return +end + +""" + wait(evt::cl.AbstractEvent) + wait(evts::Vector{cl.AbstractEvent}) + +Wait for the commands associated with the event(s) to complete, throwing a `CLError` if a +command was terminated abnormally. + +While waiting, the current task yields to others. Waiting is interruptible, in which case +the commands keep executing. +""" +function Base.wait(evt::AbstractEvent) + wait_events(AbstractEvent[evt]) + return evt +end + +function Base.wait(evts::Vector{AbstractEvent}) + isempty(evts) || wait_events(evts) + return evts +end diff --git a/src/util.jl b/src/util.jl index 19e69596..de284455 100644 --- a/src/util.jl +++ b/src/util.jl @@ -94,6 +94,7 @@ function versioninfo(io::IO=stdout) prefs = [ "default_memory_backend" => load_preference(OpenCL, "default_memory_backend"), "llvm_to_spirv_backend" => load_preference(OpenCL, "llvm_to_spirv_backend"), + "nonblocking_synchronization" => load_preference(OpenCL, "nonblocking_synchronization"), ] if any(x->!isnothing(x[2]), prefs) println(io, "Preferences:") diff --git a/test/synchronization.jl b/test/synchronization.jl index 2b904194..9c17d45e 100644 --- a/test/synchronization.jl +++ b/test/synchronization.jl @@ -71,3 +71,65 @@ end @test !Base.isexported(OpenCL, :synchronize) @test !Base.isexported(OpenCL, Symbol("@sync")) end + +# a marker that only completes once `gate` is completed. waiting for it deadlocks unless +# waiting yields to the task that completes the gate. +# +# these markers go on a queue of their own: freeing memory waits for the queue it was used +# on, blocking the thread when done by a finalizer, which would deadlock if that queue were +# waiting for a gate that only another task on this thread can complete. +const gated_queue = cl.CmdQueue() +function gated_marker() + gate = cl.UserEvent() + return gate, cl.enqueue_marker_with_wait_list(cl.AbstractEvent[gate]; queue=gated_queue) +end + +# complete `gate` from another task, after a delay that makes waiting go past polling +open_later(gate) = @async (sleep(0.1); cl.complete(gate)) + +@testset "cooperative waiting" begin + gate, evt = gated_marker() + opener = open_later(gate) + @test wait(evt) === evt + @test istaskdone(opener) + @test evt.status == :complete + + gate, evt = gated_marker() + opener = open_later(gate) + @test wait(cl.AbstractEvent[evt]) == [evt] + @test istaskdone(opener) +end + +@testset "failed commands" begin + gate, evt = gated_marker() + @async (sleep(0.1); cl.clSetUserEventStatus(gate, cl.CL_INVALID_VALUE)) + @test_throws cl.CLError wait(evt) +end + +# cancel a task waiting for the device, which is what ^C does on Julia 1.14+. (earlier +# versions throw an `InterruptException` into the task, but that cannot be done reliably +# from Julia code, as it is incorrect to `schedule` a task that has already started.) +if isdefined(Base, :CANCEL_TOKEN) +@testset "cancellation" begin + # like `Threads.@spawn`, running `f` under a cancellation token that can be cancelled + function cancellable_task(f) + src = Base.CancellationTokenSource() + task = Base.ScopedValues.with(() -> Threads.@spawn(f()), + Base.CANCEL_TOKEN => Base.CancellationToken(src)) + return task, () -> Base.cancel!(src) + end + + gate, evt = gated_marker() + task, cancel = cancellable_task(() -> wait(evt)) + sleep(0.1) + @test !istaskdone(task) + cancel() + @test_throws TaskFailedException wait(task) + @test evt.status != :complete + + # the command keeps executing, and can be waited for again + opener = open_later(gate) + @test wait(evt) === evt + @test istaskdone(opener) +end +end From 48935c35014af08d5aa71929b1d6eb2e2753b340 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 21:23:32 +0200 Subject: [PATCH 3/7] Synchronize queues cooperatively `cl.finish`, and with it `OpenCL.synchronize`, `OpenCL.@sync` and `KernelAbstractions.synchronize`, blocked the calling thread in `clFinish`. KernelAbstractions requires `synchronize` to yield instead, e.g., to overlap MPI communication with computation, and for `KA.@spawn` tasks to run concurrently. OpenCL cannot query whether a queue is idle, so wait cooperatively for a marker instead, which completes once all previously submitted commands have. The trailing `clFinish` is gone too: other tasks may submit more work to the queue while we wait, and this is what `synchronize` promises anyway. When checking for device-side exceptions, wait for the queue before taking the mailbox lock, so that other tasks can keep launching kernels in the meantime. Under the lock, the pending queues still need to be waited for before the mailbox can be read. That includes this queue, but only if a kernel was launched since we waited for it, which the launch counter (now atomic, and only published once a launch has been submitted) tells us. Mailboxes that need to be mapped for host access now do so on a queue of their own: mapping on the synchronized queue would also wait, blocking the thread, for any work other tasks submitted to it in the meantime. --- lib/cl/synchronization.jl | 13 ++++++++ src/OpenCLKernels.jl | 1 - src/compiler/exceptions.jl | 68 ++++++++++++++++++++++---------------- src/util.jl | 5 ++- test/synchronization.jl | 6 ++++ 5 files changed, 62 insertions(+), 31 deletions(-) diff --git a/lib/cl/synchronization.jl b/lib/cl/synchronization.jl index 2ea032cc..428bfaed 100644 --- a/lib/cl/synchronization.jl +++ b/lib/cl/synchronization.jl @@ -78,3 +78,16 @@ function Base.wait(evts::Vector{AbstractEvent}) isempty(evts) || wait_events(evts) return evts end + +# wait for the commands submitted to `queue` so far to complete. this does not check for +# device-side exceptions, as `finish` does. +function wait_idle(queue::CmdQueue; cancellable::Bool=true) + if nonblocking_synchronization + # there is no way to query whether a queue is idle, so wait for a marker instead + marker = enqueue_marker_with_wait_list(AbstractEvent[]; queue) + wait_events([marker]; cancellable) + else + clFinish(queue) + end + return +end diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index a7053712..98220c9f 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -42,7 +42,6 @@ end KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing KA.get_backend(::CLArray) = OpenCLBackend() -# TODO should be non-blocking KA.synchronize(::OpenCLBackend) = OpenCL.synchronize() KA.supports_float64(::OpenCLBackend) = in("cl_khr_fp64", cl.device().extensions) diff --git a/src/compiler/exceptions.jl b/src/compiler/exceptions.jl index b244c2d5..ba3f4344 100644 --- a/src/compiler/exceptions.jl +++ b/src/compiler/exceptions.jl @@ -69,14 +69,18 @@ mutable struct ExceptionMailbox # devices don't support any (in which case kernels get a zero address) const mem::Union{Nothing, cl.AbstractMemory} const address::UInt64 - # whether host access requires mapping the memory (coarse-grained SVM, buffers) + # whether host access requires mapping the memory (coarse-grained SVM, buffers), which + # happens on a queue of its own: mapping on a queue that is being synchronized would + # also wait for work that other tasks submitted to it since. const mapped::Bool + const queue::Union{Nothing, cl.CmdQueue} # serialize launches and checks for this mailbox const lock::ReentrantLock # queues that have launched kernels since the mailbox was last checked on them const pending::Set{cl.CmdQueue} - # assigned under the lock; distinguishes work-items from different launches - launch_id::UInt64 + # the last launch, assigned under the lock after submission; distinguishes work-items + # from different launches + Base.@atomic launch_id::UInt64 end # Device-scoped atomics cannot protect a mailbox shared by different devices. These strong @@ -90,8 +94,7 @@ const exception_mailbox_backends = (:usm, :svm_fine, :svm_coarse, :buffer) # allocate a mailbox in memory all devices in the context can reach by pointer, trying the # given backends in order. -function allocate_exception_mailbox(ctx::cl.Context, queue::cl.CmdQueue=cl.queue(); - backends=exception_mailbox_backends) +function allocate_exception_mailbox(ctx::cl.Context; backends=exception_mailbox_backends) devs = ctx.devices sz = sizeof(ExceptionInfo_st) mem, mapped = cl.context!(ctx) do @@ -123,22 +126,24 @@ function allocate_exception_mailbox(ctx::cl.Context, queue::cl.CmdQueue=cl.queue @warn """Device-side exceptions cannot be reported on $(join(map(dev -> dev.name, devs), ", ")): \ the device does not support any host-accessible memory to put the exception mailbox in. Kernels that throw will complete silently.""" - return ExceptionMailbox(nothing, UInt64(0), false, ReentrantLock(), + return ExceptionMailbox(nothing, UInt64(0), false, nothing, ReentrantLock(), Set{cl.CmdQueue}(), UInt64(0)) end - mailbox = ExceptionMailbox(mem, UInt64(UInt(pointer(mem))), mapped, ReentrantLock(), - Set{cl.CmdQueue}(), UInt64(0)) - with_exception_mailbox(mailbox, queue) do ptr + queue = mapped ? cl.context!(cl.CmdQueue, ctx) : nothing + mailbox = ExceptionMailbox(mem, UInt64(UInt(pointer(mem))), mapped, queue, + ReentrantLock(), Set{cl.CmdQueue}(), UInt64(0)) + with_exception_mailbox(mailbox) do ptr unsafe_store!(ptr, ExceptionInfo_st()) end return mailbox end -# run `f` with a host pointer to the mailbox, mapping its memory through `queue` as needed -function with_exception_mailbox(f, mailbox::ExceptionMailbox, queue::cl.CmdQueue) +# run `f` with a host pointer to the mailbox, mapping its memory as needed +function with_exception_mailbox(f, mailbox::ExceptionMailbox) sz = sizeof(ExceptionInfo_st) mem = mailbox.mem + queue = mailbox.queue if mem isa cl.Buffer ptr, _ = cl.enqueue_map(mem, sz, :rw; queue, blocking=true) try @@ -174,15 +179,18 @@ function launch_with_exception_mailbox(kernel::cl.Kernel, args...; ctx, dev, queue = cl.context(), cl.device(), cl.queue() mailbox = Base.@lock exception_mailboxes_lock begin get!(exception_mailboxes, (ctx, dev)) do - allocate_exception_mailbox(ctx, queue) + allocate_exception_mailbox(ctx) end end Base.@lock mailbox.lock begin mailbox.mem === nothing || push!(indirect_memory, mailbox.mem) - mailbox.launch_id += UInt64(1) + launch_id = (Base.@atomic mailbox.launch_id) + UInt64(1) state = KernelState(rng_state ? Base.rand(UInt32) : UInt32(0), mailbox.address, - mailbox.launch_id) + launch_id) result = cl.call(kernel, state, args...; indirect_memory, rng_state, kwargs...) + # only publish the launch once it has been submitted, so that `check_exceptions` + # can tell whether kernels have been submitted while it waited without the lock + Base.@atomic mailbox.launch_id = launch_id if mailbox.mem !== nothing push!(mailbox.pending, queue) end @@ -201,34 +209,36 @@ The exception mailbox is shared by queues targeting the same device in a context may wait for and surface an exception from another queue on that device. """ function check_exceptions(queue::cl.CmdQueue; rethrow::Bool=true) - # Finalizers cannot yield while waiting for a contended Julia lock. Finish the queue - # without touching mailbox bookkeeping; an ordinary check will finish it again and - # consume the report. This also keeps finalizers independent of task-local state. + # Finalizers cannot yield while waiting for a contended Julia lock. Wait for the queue + # without touching mailbox bookkeeping; an ordinary check will wait again and consume + # the report. This also keeps finalizers independent of task-local state. if !rethrow - cl.clFinish(queue) + cl.wait_idle(queue; cancellable=false) return end + mailbox = Base.@lock exception_mailboxes_lock begin get(exception_mailboxes, (queue.context, queue.device), nothing) end - if mailbox === nothing - cl.clFinish(queue) - return - end + + # Wait before taking the mailbox lock, so that other tasks can keep launching kernels. + launch_id = mailbox === nothing ? nothing : Base.@atomic mailbox.launch_id + cl.wait_idle(queue) + mailbox === nothing && return + Base.@lock mailbox.lock begin - # Keep this raw: `cl.finish` delegates here in order to hold the mailbox lock across - # both synchronization and inspection. - cl.clFinish(queue) isempty(mailbox.pending) && return # The mailbox cannot be mapped or read while another queue may still access it. - # Finish all queues that launched with its address; the launch-side lock prevents a - # new enqueue from appearing between this loop and the reset below. + # Wait for all queues that launched with its address, which only excludes this + # queue if nothing was launched since waiting for it above; the launch-side lock + # prevents a new enqueue from appearing between this loop and the reset below. for pending in mailbox.pending - pending == queue || cl.clFinish(pending) + pending == queue && (Base.@atomic mailbox.launch_id) == launch_id && continue + cl.wait_idle(pending) end - exc = with_exception_mailbox(mailbox, queue) do ptr + exc = with_exception_mailbox(mailbox) do ptr status_ptr = convert(Ptr{Int32}, ptr) unsafe_load(status_ptr) == 0 && return nothing info = unsafe_load(ptr) diff --git a/src/util.jl b/src/util.jl index de284455..65eff264 100644 --- a/src/util.jl +++ b/src/util.jl @@ -188,6 +188,9 @@ Wait for all operations submitted to `queue` to finish. By default, this is the queue of the current task, `cl.queue()`; operations submitted by other tasks, which use their own queues, are not waited for. +While waiting, the current task yields, so other tasks can run on its thread. Waiting can be +interrupted, in which case the device keeps executing the operations. + If a kernel threw an exception, it is reported here as a [`KernelException`](@ref). Because queues targeting the same device share their exception state, this may also wait for, and report a failure from, another queue on that device. @@ -205,7 +208,7 @@ end Run expression `ex`, then wait for the operations it submitted to the current task's queue to finish, using [`OpenCL.synchronize`](@ref). Returns the value of `ex`. -This blocks the calling thread while waiting. Only the current task's queue is waited for: +Only the current task's queue is waited for: work that `ex` submits to other queues, e.g., inside a `cl.queue!(q) do ... end` block, or from Julia tasks it spawns (unlike `Base.@sync`), needs to be synchronized separately. """ diff --git a/test/synchronization.jl b/test/synchronization.jl index 9c17d45e..13535121 100644 --- a/test/synchronization.jl +++ b/test/synchronization.jl @@ -98,6 +98,12 @@ open_later(gate) = @async (sleep(0.1); cl.complete(gate)) opener = open_later(gate) @test wait(cl.AbstractEvent[evt]) == [evt] @test istaskdone(opener) + + gate, evt = gated_marker() + opener = open_later(gate) + OpenCL.synchronize(gated_queue) + @test istaskdone(opener) + @test evt.status == :complete end @testset "failed commands" begin From b385ef7e58f57be2c25c13d1203f571865010167 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 21:23:32 +0200 Subject: [PATCH 4/7] Wait for blocking transfers cooperatively Blocking reads, writes and copies (e.g., `Array(x)`, or `copyto!` to host memory) passed `blocking=true` to the driver, which then blocked the calling thread until all earlier work on the queue had completed too. Submit these commands without blocking in the driver, and wait for their event cooperatively instead. That wait cannot be interrupted: the command accesses memory that the caller may release as soon as we return. This also makes `enqueue_copy` between buffers honor `blocking`, which it ignored before. Maps still block in the driver, as they are only issued after synchronizing. --- lib/cl/event.jl | 21 +++++++++++++++------ lib/cl/memory/buffer.jl | 10 +++++----- lib/cl/memory/svm.jl | 4 ++-- lib/cl/memory/usm.jl | 4 ++-- test/synchronization.jl | 29 ++++++++++++++++++++++++++++- 5 files changed, 52 insertions(+), 16 deletions(-) diff --git a/lib/cl/event.jl b/lib/cl/event.jl index e0462959..6c2f6914 100644 --- a/lib/cl/event.jl +++ b/lib/cl/event.jl @@ -28,27 +28,36 @@ end NannyEvent(evt::Event, obj; retain=false) = NannyEvent(evt.id, obj; retain) -macro return_event(evt) +# commands that should block are submitted without blocking in the driver, and waited for +# cooperatively instead (see `wait_events`). that wait cannot be interrupted, as the caller +# may release the memory the command accesses once we return. +macro return_event(evt, blocking=false) quote evt = $(esc(evt)) - try - return Event(evt, retain=false) + evt = try + Event(evt, retain=false) catch err + $(esc(blocking)) && unchecked_clWaitForEvents(1, Ref(evt)) clReleaseEvent(evt) throw(err) end + $(esc(blocking)) && wait_events([evt]; cancellable=false) + return evt end end -macro return_nanny_event(evt, obj) +macro return_nanny_event(evt, obj, blocking=false) quote evt = $(esc(evt)) - try - return NannyEvent(evt, $(esc(obj))) + evt = try + NannyEvent(evt, $(esc(obj))) catch err + $(esc(blocking)) && unchecked_clWaitForEvents(1, Ref(evt)) clReleaseEvent(evt) throw(err) end + $(esc(blocking)) && wait_events([evt]; cancellable=false) + return evt end end diff --git a/lib/cl/memory/buffer.jl b/lib/cl/memory/buffer.jl index 974f6f55..106f2386 100644 --- a/lib/cl/memory/buffer.jl +++ b/lib/cl/memory/buffer.jl @@ -89,9 +89,9 @@ function enqueue_read(dst::Ptr, src::Buffer, src_off::Int, nbytes::Int; evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for] GC.@preserve wait_for begin ret_evt = Ref{cl_event}() - clEnqueueReadBuffer(queue(), src, blocking, src_off, nbytes, dst, + clEnqueueReadBuffer(queue(), src, false, src_off, nbytes, dst, n_evts, evt_ids, ret_evt) - @return_nanny_event(ret_evt[], dst) + @return_nanny_event(ret_evt[], dst, blocking) end end enqueue_read(dst::Ptr, src::Buffer, nbytes; kwargs...) = @@ -104,9 +104,9 @@ function enqueue_write(dst::Buffer, dst_off::Int, src::Ptr, nbytes::Int; evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for] GC.@preserve wait_for begin ret_evt = Ref{cl_event}() - clEnqueueWriteBuffer(queue(), dst, blocking, dst_off, nbytes, src, + clEnqueueWriteBuffer(queue(), dst, false, dst_off, nbytes, src, n_evts, evt_ids, ret_evt) - @return_nanny_event(ret_evt[], dst) + @return_nanny_event(ret_evt[], dst, blocking) end end enqueue_write(dst::Buffer, src::Ptr, nbytes; kwargs...) = @@ -122,7 +122,7 @@ function enqueue_copy(dst::Buffer, dst_off::Int, src::Buffer, src_off::Int, ret_evt = Ref{cl_event}() clEnqueueCopyBuffer(queue(), src, dst, src_off, dst_off, nbytes, n_evts, evt_ids, ret_evt) - @return_event ret_evt[] + @return_event ret_evt[] blocking end end enqueue_copy(dst::Buffer, src::Buffer, N; kwargs...) = diff --git a/lib/cl/memory/svm.jl b/lib/cl/memory/svm.jl index 837062b6..c64a1c65 100644 --- a/lib/cl/memory/svm.jl +++ b/lib/cl/memory/svm.jl @@ -86,8 +86,8 @@ function enqueue_svm_copy( evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for] return GC.@preserve wait_for begin ret_evt = Ref{cl_event}() - clEnqueueSVMMemcpy(queue, blocking, dst, src, nbytes, n_evts, evt_ids, ret_evt) - @return_event ret_evt[] + clEnqueueSVMMemcpy(queue, false, dst, src, nbytes, n_evts, evt_ids, ret_evt) + @return_event ret_evt[] blocking end end diff --git a/lib/cl/memory/usm.jl b/lib/cl/memory/usm.jl index 819a7c21..281cd84c 100644 --- a/lib/cl/memory/usm.jl +++ b/lib/cl/memory/usm.jl @@ -208,8 +208,8 @@ function enqueue_usm_copy( return GC.@preserve wait_for begin com = (CL_NULL, C_NULL) ret_evt = (dst in com || src in com) ? C_NULL : Ref{cl_event}() - clEnqueueMemcpyINTEL(queue, blocking, dst, src, nbytes, n_evts, evt_ids, ret_evt) - @return_event ret_evt[] + clEnqueueMemcpyINTEL(queue, false, dst, src, nbytes, n_evts, evt_ids, ret_evt) + @return_event ret_evt[] blocking end end diff --git a/test/synchronization.jl b/test/synchronization.jl index 13535121..347a4ed7 100644 --- a/test/synchronization.jl +++ b/test/synchronization.jl @@ -77,8 +77,10 @@ end # # these markers go on a queue of their own: freeing memory waits for the queue it was used # on, blocking the thread when done by a finalizer, which would deadlock if that queue were -# waiting for a gate that only another task on this thread can complete. +# waiting for a gate that only another task on this thread can complete. for the same +# reason, the only memory used on that queue is kept alive. const gated_queue = cl.CmdQueue() +const gated_array = cl.queue!(() -> OpenCL.ones(Float32, 4), gated_queue) function gated_marker() gate = cl.UserEvent() return gate, cl.enqueue_marker_with_wait_list(cl.AbstractEvent[gate]; queue=gated_queue) @@ -104,6 +106,12 @@ open_later(gate) = @async (sleep(0.1); cl.complete(gate)) OpenCL.synchronize(gated_queue) @test istaskdone(opener) @test evt.status == :complete + + # blocking copies + gate, evt = gated_marker() + opener = open_later(gate) + @test cl.queue!(() -> Array(gated_array), gated_queue) == ones(Float32, 4) + @test istaskdone(opener) end @testset "failed commands" begin @@ -137,5 +145,24 @@ if isdefined(Base, :CANCEL_TOKEN) opener = open_later(gate) @test wait(evt) === evt @test istaskdone(opener) + + # blocking transfers only return once they have completed, even when cancelled, as + # the caller may release the memory they access + gate, evt = gated_marker() + dev = cl.device() + task, cancel = cancellable_task() do + cl.device!(dev) + cl.queue!(() -> Array(gated_array), gated_queue) + end + sleep(0.1) + cancel() + sleep(0.1) + @test !istaskdone(task) + cl.complete(gate) + try + wait(task) + catch + end + @test istaskdone(task) end end From dc8dd3f69b5aebd946a92eb45730c803a92e614f Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 00:00:08 +0200 Subject: [PATCH 5/7] Wait for commands on CPU devices using completion callbacks MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Handing the blocking wait to a worker thread works well for GPUs, but slows down devices that execute on the host's CPU cores (e.g., PoCL, or Intel's CPU runtime): waking the worker as the commands start, and polling their status for a long time, compete with them for those cores. On 4 cores with PoCL, kernels of 0.2-1 ms took 25-35% longer than with a blocking wait. For such devices, only poll for at most 10 µs, and then have the driver notify us when the commands complete, using `clSetEventCallback`, which GPUToolbox 3.3 supports. Kernels then take about as long as with a blocking wait. GPUs keep using worker threads: NVIDIA's driver delivers a few percent of these callbacks about 20 ms late. --- Project.toml | 2 +- lib/cl/api.jl | 2 +- lib/cl/synchronization.jl | 63 ++++++++++++++++++++++++++++++++------- 3 files changed, 55 insertions(+), 12 deletions(-) diff --git a/Project.toml b/Project.toml index 9ab45ad4..af34fed0 100644 --- a/Project.toml +++ b/Project.toml @@ -33,7 +33,7 @@ SPIRVIntrinsics = {path = "lib/intrinsics"} Adapt = "4" GPUArrays = "11.2.1" GPUCompiler = "2.9" -GPUToolbox = "3.2" +GPUToolbox = "3.3" KernelAbstractions = "0.9.38" LLVM = "9.6" LinearAlgebra = "1" diff --git a/lib/cl/api.jl b/lib/cl/api.jl index 86721ca9..7fe32697 100644 --- a/lib/cl/api.jl +++ b/lib/cl/api.jl @@ -1,5 +1,5 @@ import OpenCL_jll -using GPUToolbox: @gcsafe_ccall, cooperative_wait +using GPUToolbox: GPUToolbox, @gcsafe_ccall, cooperative_wait const libopencl = OpenCL_jll.libopencl diff --git a/lib/cl/synchronization.jl b/lib/cl/synchronization.jl index 428bfaed..983148a7 100644 --- a/lib/cl/synchronization.jl +++ b/lib/cl/synchronization.jl @@ -2,8 +2,14 @@ # # Waiting for the device should not block the calling thread in the OpenCL driver, but # yield to the Julia scheduler so that other tasks can run in the meantime. GPUToolbox's -# `cooperative_wait` implements this by polling the status of commands for a while, and -# then handing the blocking wait to a worker thread. +# `cooperative_wait` implements this, in one of two ways: +# - for GPUs, by polling the status of commands for a while, and then handing a blocking +# wait to a worker thread; +# - for devices that execute on the host's CPU cores, by polling only briefly, and then +# having the driver notify us when commands complete. waking a worker, or polling for +# longer, would compete with the commands for those cores, slowing them down +# considerably. (drivers for GPUs may deliver these notifications late, NVIDIA's by +# about 20 ms.) # whether to wait for the device cooperatively. disable to block in the OpenCL driver # instead, e.g., for bisecting issues or comparing against the blocking behavior. @@ -21,14 +27,44 @@ function iscomplete(evt::AbstractEvent) end # commands only need to start executing once the queue they were submitted to has been -# flushed. waiting in the driver flushes implicitly, but polling does not. +# flushed. waiting in the driver flushes implicitly, but polling does not. returns the +# queue, or `C_NULL` for user events, which do not belong to one. function flush_queue(evt::AbstractEvent) queue = Ref{cl_command_queue}() clGetEventInfo(evt, CL_EVENT_COMMAND_QUEUE, sizeof(cl_command_queue), queue, C_NULL) - # user events do not belong to a queue queue[] == C_NULL || clFlush(queue[]) + return queue[] +end + +# whether commands on `queue` execute on the host's CPU cores +const cpu_devices = Dict{cl_device_id, Bool}() +const cpu_devices_lock = ReentrantLock() +function executes_on_cpu(queue::cl_command_queue) + queue == C_NULL && return false + device = Ref{cl_device_id}() + clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), device, C_NULL) + return @lock cpu_devices_lock get!(cpu_devices, device[]) do + type = Ref{cl_device_type}() + clGetDeviceInfo(device[], CL_DEVICE_TYPE, sizeof(cl_device_type), type, C_NULL) + type[] & CL_DEVICE_TYPE_CPU != 0 + end +end + +# driver notification that a command has completed +function notify_completion(::cl_event, ::Cint, payload::Ptr{Cvoid}) + GPUToolbox.signal_completion(payload) return end +subscribe_completion(evt, payload) = + clSetEventCallback(evt, CL_COMPLETE, + @cfunction(notify_completion, Cvoid, (cl_event, Cint, Ptr{Cvoid})), + payload) + +# block until the events have completed, without checking for errors +blocking_wait(evts::Vector{<:AbstractEvent}) = + GC.@preserve evts unchecked_clWaitForEvents(length(evts), event_ids(evts)) +blocking_wait(evt::AbstractEvent) = + GC.@preserve evt unchecked_clWaitForEvents(1, Ref(pointer(evt))) # wait for events to complete, throwing a `CLError` if a command was terminated abnormally. # @@ -39,17 +75,24 @@ end function wait_events(evts::Vector{<:AbstractEvent}; cancellable::Bool=true) if nonblocking_synchronization try - foreach(flush_queue, evts) - cooperative_wait(evts; isdone=evts -> all(iscomplete, evts), cancellable) do evts - GC.@preserve evts unchecked_clWaitForEvents(length(evts), event_ids(evts)) + on_cpu = true + for evt in evts + on_cpu &= executes_on_cpu(flush_queue(evt)) + end + if on_cpu + for evt in evts + cooperative_wait(blocking_wait, evt; subscribe=subscribe_completion, + isdone=iscomplete, spin=10e-6, cancellable) + end + else + cooperative_wait(blocking_wait, evts; isdone=evts -> all(iscomplete, evts), + cancellable) end catch # waits that cannot be interrupted only throw once the commands have completed # (or waiting failed, in which case this blocks), but host memory still needs to # be synchronized before the caller unwinds - cancellable || GC.@preserve evts begin - unchecked_clWaitForEvents(length(evts), event_ids(evts)) - end + cancellable || blocking_wait(evts) rethrow() end end From e776cf66f8fe1ea1b41825df0457d082b0480b89 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 12:40:28 +0200 Subject: [PATCH 6/7] Require GPUToolbox 3.3.1 On GPUs, waits are handed to GPUToolbox's worker threads, which 3.3.1 keeps from deadlocking: workers no longer run finalizers, which could block (e.g., freeing memory waits for the device) before the worker notifies its waiter (JuliaGPU/GPUToolbox.jl#26). --- Project.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index af34fed0..55d81412 100644 --- a/Project.toml +++ b/Project.toml @@ -33,7 +33,7 @@ SPIRVIntrinsics = {path = "lib/intrinsics"} Adapt = "4" GPUArrays = "11.2.1" GPUCompiler = "2.9" -GPUToolbox = "3.3" +GPUToolbox = "3.3.1" KernelAbstractions = "0.9.38" LLVM = "9.6" LinearAlgebra = "1" From 3f57841782bf7ffa76b74dfa3d74729ca272afa0 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 15:03:39 +0200 Subject: [PATCH 7/7] Require GPUToolbox 3.3.2 On CPU devices, which wait for completion callbacks, 3.3.2 keeps the woken task from spinning on a lock held by the driver's callback thread after preempting it, which made longer PoCL kernels slower when all cores are busy (JuliaGPU/GPUToolbox.jl#27). --- Project.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index 55d81412..cbaf7cdf 100644 --- a/Project.toml +++ b/Project.toml @@ -33,7 +33,7 @@ SPIRVIntrinsics = {path = "lib/intrinsics"} Adapt = "4" GPUArrays = "11.2.1" GPUCompiler = "2.9" -GPUToolbox = "3.3.1" +GPUToolbox = "3.3.2" KernelAbstractions = "0.9.38" LLVM = "9.6" LinearAlgebra = "1"