Implement queue

Author: vchuravy

src/pocl/backend.jl

@@ -110,6 +110,7 @@ function (obj::KA.Kernel{POCLBackend})(args...; ndrange=nothing, workgroupsize=n
     local_size = items
     kernel(ctx, args...; global_size, local_size)
 
+    cl.finish(cl.queue()) # TODO, would waiting on an event be cheaper?
     return nothing
 end
 

src/pocl/nanoOpenCL.jl

@@ -476,6 +476,10 @@ const cl_kernel_work_group_info = cl_uint
 
 const cl_kernel_sub_group_info = cl_uint
 
+const cl_device_svm_capabilities = cl_bitfield
+
+const cl_command_queue_properties = cl_bitfield
+
 @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
@@ -589,6 +593,20 @@ end
                                             event::Ptr{cl_event})::cl_int
 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
+end
+
+@checked function clReleaseCommandQueue(command_queue)
+    @ccall libopencl.clReleaseCommandQueue(command_queue::cl_command_queue)::cl_int
+end
+
+@checked function clFinish(command_queue)
+    @ccall libopencl.clFinish(command_queue::cl_command_queue)::cl_int
+end
+
 # Init
 
 # lazy initialization
@@ -886,16 +904,6 @@ function Context(device::Device)
     return Context(ctx_id)
 end
 
-mutable struct Event
-    const id::cl_event
-
-    function Event(evt_id)
-        evt = new(evt_id)
-        finalizer(clReleaseEvent, evt)
-        return evt
-    end
-end
-
 mutable struct Program
     const id::cl_program
 
@@ -1138,7 +1146,7 @@ function set_args!(k::Kernel, args...)
 end
 
 function enqueue_kernel(k::Kernel, global_work_size, local_work_size=nothing;
-                        global_work_offset=nothing, wait_on::Vector{Event}=Event[])
+                        global_work_offset=nothing)
     max_work_dim = device().max_work_item_dims
     work_dim     = length(global_work_size)
     if work_dim > max_work_dim
@@ -1181,29 +1189,24 @@ function enqueue_kernel(k::Kernel, global_work_size, local_work_size=nothing;
         # null local size means OpenCL decides
     end
 
-    if !isempty(wait_on)
-        n_events = cl_uint(length(wait_on))
-        wait_event_ids = [evt.id for evt in wait_on]
-    else
-        n_events = cl_uint(0)
-        wait_event_ids = C_NULL
-    end
+    n_events = cl_uint(0)
+    wait_event_ids = C_NULL
+    ret_event = C_NULL
 
-    ret_event = Ref{cl_event}()
     clEnqueueNDRangeKernel(queue(), k, cl_uint(work_dim), goffset, gsize, lsize,
                            n_events, wait_event_ids, ret_event)
-    return Event(ret_event[], retain=false)
+    return nothing
 end
 
 function call(k::Kernel, args...; global_size=(1,), local_size=nothing,
-              global_work_offset=nothing, wait_on::Vector{Event}=Event[],
+              global_work_offset=nothing,
               svm_pointers::Vector{Ptr{Cvoid}}=Ptr{Cvoid}[])
     set_args!(k, args...)
     if !isempty(svm_pointers)
         clSetKernelExecInfo(k, CL_KERNEL_EXEC_INFO_SVM_PTRS,
                             sizeof(svm_pointers), svm_pointers)
     end
-    enqueue_kernel(k, global_size, local_size; global_work_offset, wait_on)
+    enqueue_kernel(k, global_size, local_size; global_work_offset)
 end
 
 # convert the argument values to match the kernel's signature (specified by the user)
@@ -1272,4 +1275,36 @@ function Base.getproperty(ki::KernelWorkGroupInfo, s::Symbol)
     end
 end
 
+mutable struct CmdQueue
+    const id::cl_command_queue
+
+    function CmdQueue(q_id::cl_command_queue)
+        q = new(q_id)
+        finalizer(q) do _
+          clReleaseCommandQueue(q)
+        end
+        return q
+    end
+end
+
+Base.unsafe_convert(::Type{cl_command_queue}, q::CmdQueue) = q.id
+
+function CmdQueue()
+    flags=cl_command_queue_properties(0)
+    err_code = Ref{Cint}()
+    queue_id = clCreateCommandQueue(context(), device(), flags, err_code)
+    if err_code[] != CL_SUCCESS
+        if queue_id != C_NULL
+            clReleaseCommandQueue(queue_id)
+        end
+        throw(CLError(err_code[]))
+    end
+    return CmdQueue(queue_id)
+end
+
+function finish(q::CmdQueue)
+    clFinish(q)
+    return q
+end
+
 end

src/pocl/pocl.jl

@@ -27,6 +27,7 @@ function device()
     end::cl.Device
 end
 
+# TODO: add a device context dict
 function context()
     get!(task_local_storage(), :POCLContext) do
         cl.Context(device())
@@ -36,7 +37,7 @@ end
 function queue()
     get!(task_local_storage(), :POCLQueue) do
         cl.CmdQueue()
-    end::CmdQueue
+    end::cl.CmdQueue
 end
 
 using GPUCompiler