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/Project.toml b/Project.toml index 0236d89d..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.1" +GPUToolbox = "3.3.2" KernelAbstractions = "0.9.38" LLVM = "9.6" LinearAlgebra = "1" 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 533a1add..7fe32697 100644 --- a/lib/cl/api.jl +++ b/lib/cl/api.jl @@ -1,4 +1,5 @@ import OpenCL_jll +using GPUToolbox: GPUToolbox, @gcsafe_ccall, cooperative_wait 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/event.jl b/lib/cl/event.jl index 67a53010..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 @@ -143,26 +152,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/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/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/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/lib/cl/synchronization.jl b/lib/cl/synchronization.jl new file mode 100644 index 00000000..983148a7 --- /dev/null +++ b/lib/cl/synchronization.jl @@ -0,0 +1,136 @@ +# 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, 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. +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. 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) + 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. +# +# 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 + 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 || blocking_wait(evts) + 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 + +# 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/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 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 19e69596..65eff264 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:") @@ -187,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. @@ -204,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 2b904194..347a4ed7 100644 --- a/test/synchronization.jl +++ b/test/synchronization.jl @@ -71,3 +71,98 @@ 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. 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) +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) + + gate, evt = gated_marker() + opener = open_later(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 + 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) + + # 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