Update: implicit event reference counting.

Uses overloaded assignment and finalisation routine
to call OpenCL reference counting routines for the
underlying event objects.
This is needed for when library users store their
own copies of event objects for dependency management,

src/Focal.f90

@@ -92,6 +92,8 @@ module Focal
     type :: fclEvent
     !! Type wrapper for OpenCL event pointers
     integer(c_intptr_t) :: cl_event = -1             !! OpenCL event pointer
+    contains
+      final :: fclReleaseEvent                       !! Decrement cl reference counter
   end type fclEvent
 
   type :: fclCommandQ
@@ -1395,6 +1397,31 @@ end subroutine fclWaitEventList
 
   end interface fclWait
 
+  interface assignment(=)
+
+    module subroutine fclEventCopy(target,source)
+      !! Overloaded assignment for event assignment.
+      !!  Handles opencl reference counting for the underlying event object
+      type(fclEvent), intent(inout) :: target
+      type(fclEvent), intent(in) :: source
+    end subroutine fclEventCopy
+
+  end interface
+
+  interface 
+
+    module subroutine fclReleaseEvent(event)
+      !! Light weight wrapper for clReleaseEvent (decrement reference count)
+      type(fclEvent), intent(in) :: event               !! Focal event object to release
+    end subroutine fclReleaseEvent
+
+    module subroutine fclRetainEvent(event)
+      !! Light weight wrapper for clRetainEvent (increment reference count)
+      type(fclEvent), intent(in) :: event               !! Focal event object to retain
+    end subroutine fclRetainEvent
+
+  end interface
+
   interface fclSetDependency
     !! Generic interface to set pre-requisite events for the next enqueued action.
     !!  This does not append to any existing dependencies - it overwrites the dependency list.

src/Focal_Memory.f90

@@ -331,25 +331,21 @@
   module procedure fclMemWriteScalar !(memObject,hostBufferPtr,nBytesPattern)
 
     integer(c_int32_t) :: errcode
+    type(fclEvent), target :: writeEvent
 
     call fclDbgCheckBufferInit(memObject,'fclMemWriteScalar')
 
-    ! Decrement event reference counter
-    if (memObject%cmdq%lastWriteEvent%cl_event > 0) then
-      errcode = clReleaseEvent(memObject%cmdq%lastWriteEvent%cl_event)
-      call fclErrorHandler(errcode,'fclMemWriteScalar','clReleaseEvent') 
-    end if
-
     errcode = clEnqueueFillBuffer(memObject%cmdq%cl_command_queue, &
                 memObject%cl_mem, hostBufferPtr, nBytesPattern, &
                 int(0,c_size_t), memObject%nBytes, &
                 memObject%cmdq%nDependency, memObject%cmdq%dependencyListPtr, &
-                c_loc(memObject%cmdq%lastWriteEvent%cl_event))
+                c_loc(writeEvent%cl_event))
 
     call fclPopDependencies(memObject%cmdq)
-    fclLastWriteEvent = memObject%cmdq%lastWriteEvent
+    fclLastWriteEvent = writeEvent
+    memObject%cmdq%lastWriteEvent = writeEvent
 
-    call memObject%pushProfileEvent(memObject%cmdq%lastWriteEvent,1)
+    call memObject%pushProfileEvent(writeEvent,1)
 
     call fclErrorHandler(errcode,'fclMemWriteScalar','clEnqueueFillBuffer')
 
@@ -391,6 +387,7 @@
 
     integer(c_int32_t) :: errcode
     integer(c_int32_t) :: blocking_write
+    type(fclEvent), target :: writeEvent
 
     call fclDbgCheckBufferInit(memObject,'fclMemWrite')
     call fclDbgCheckBufferSize(memObject,nBytes,'fclMemWrite')
@@ -401,21 +398,16 @@
       blocking_write = CL_FALSE
     end if
 
-    ! Decrement event reference counter
-    if (memObject%cmdq%lastWriteEvent%cl_event > 0) then
-      errcode = clReleaseEvent(memObject%cmdq%lastWriteEvent%cl_event)
-      call fclErrorHandler(errcode,'fclMemWrite','clReleaseEvent') 
-    end if
-
     errcode = clEnqueueWriteBuffer(memObject%cmdq%cl_command_queue,memObject%cl_mem, &
           blocking_write,int(0,c_size_t),nBytes,hostBufferPtr, &
           memObject%cmdq%nDependency, memObject%cmdq%dependencyListPtr, &
-          c_loc(memObject%cmdq%lastWriteEvent%cl_event))
+          c_loc(writeEvent%cl_event))
 
     call fclPopDependencies(memObject%cmdq)
-    fclLastWriteEvent = memObject%cmdq%lastWriteEvent
+    fclLastWriteEvent = writeEvent
+    memObject%cmdq%lastWriteEvent = writeEvent
 
-    call memObject%pushProfileEvent(memObject%cmdq%lastWriteEvent,1)
+    call memObject%pushProfileEvent(writeEvent,1)
 
     call fclErrorHandler(errcode,'fclMemWrite','clEnqueueWriteBuffer')
 
@@ -460,6 +452,7 @@
 
     integer(c_int32_t) :: errcode
     integer(c_int32_t) :: blocking_read
+    type(fclEvent), target :: readEvent
 
     call fclDbgCheckBufferInit(memObject,'fclMemRead')
     call fclDbgCheckBufferSize(memObject,nBytes,'fclMemRead')
@@ -470,21 +463,16 @@
       blocking_read = CL_FALSE
     end if
 
-    ! Decrement event reference counter
-    if (memObject%cmdq%lastReadEvent%cl_event > 0) then
-      errcode = clReleaseEvent(memObject%cmdq%lastReadEvent%cl_event)
-      call fclErrorHandler(errcode,'fclMemRead','clReleaseEvent') 
-    end if
-
     errcode = clEnqueueReadBuffer(memObject%cmdq%cl_command_queue,memObject%cl_mem, &
           blocking_read,int(0,c_size_t),nBytes,hostBufferPtr, &
           memObject%cmdq%nDependency, memObject%cmdq%dependencyListPtr, &
-          c_loc(memObject%cmdq%lastReadEvent%cl_event))
+          c_loc(readEvent%cl_event))
 
     call fclPopDependencies(memObject%cmdq)
-    fclLastReadEvent = memObject%cmdq%lastReadEvent
+    fclLastReadEvent = readEvent
+    memObject%cmdq%lastReadEvent = readEvent
 
-    call memObject%pushProfileEvent(memObject%cmdq%lastReadEvent,2)
+    call memObject%pushProfileEvent(readEvent,2)
 
     call fclErrorHandler(errcode,'fclMemRead','clEnqueueReadBuffer')
 
@@ -528,6 +516,7 @@
   module procedure fclMemCopy !(memObject1,memObject2)
 
     integer(c_int32_t) :: errcode
+    type(fclEvent), target :: copyEvent
 
     if (memObject2%nBytes < 0) then
       ! Source object is uninitialised: nothing to copy
@@ -559,23 +548,18 @@
 
       call fclDbgCheckCopyBufferSize(memObject1,memObject2)
 
-      ! Decrement event reference counter
-      if (memObject1%cmdq%lastCopyEvent%cl_event > 0) then
-        errcode = clReleaseEvent(memObject1%cmdq%lastCopyEvent%cl_event)
-        call fclErrorHandler(errcode,'fclMemCopy','clReleaseEvent') 
-      end if
-
       errcode = clEnqueueCopyBuffer(memObject1%cmdq%cl_command_queue, &
                 memObject2%cl_mem, memObject1%cl_mem, &
                 int(0,c_size_t), int(0,c_size_t), &
                 memObject2%nBytes, &
                 memObject1%cmdq%nDependency, memObject1%cmdq%dependencyListPtr, &
-                c_loc(memObject1%cmdq%lastCopyEvent%cl_event))
+                c_loc(copyEvent%cl_event))
 
       call fclPopDependencies(memObject1%cmdq)
-      fclLastCopyEvent = memObject1%cmdq%lastCopyEvent
+      fclLastCopyEvent = copyEvent
+      memObject1%cmdq%lastCopyEvent = copyEvent
 
-      call memObject1%pushProfileEvent(memObject1%cmdq%lastCopyEvent,3)
+      call memObject1%pushProfileEvent(copyEvent,3)
 
       call fclErrorHandler(errcode,'fclMemCopy','clEnqueueCopyBuffer')
 

src/Focal_Profile.f90

@@ -191,10 +191,6 @@
       return
     end if
 
-    ! Increment event reference counter
-    errcode = clRetainEvent(event%cl_event)
-    call fclErrorHandler(errcode,'fclSetDependencyEvent','clRetainEvent')
-
     ! Save event
     container%profileEvents(container%nProfileEvent) = event
 

src/Focal_Setup.f90

@@ -670,6 +670,7 @@
     type(fclCommandQ), pointer :: cmdQ
     type(c_ptr) :: localSizePtr
     integer :: nArg
+    type(fclEvent), target :: kernelEvent
 
     ! Check global size has been set
     if (sum(abs(kernel%global_work_size)) == 0) then
@@ -698,26 +699,21 @@
     call fclProcessKernelArgs(kernel,cmdq,narg,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,&
                                 a10,a11,a12,a13,a14,a15,a16,a17,a18,a19)
 
-    ! Decrement event reference counter
-    if (cmdQ%lastKernelEvent%cl_event > 0) then
-      errcode = clReleaseEvent(cmdQ%lastKernelEvent%cl_event)
-      call fclErrorHandler(errcode,'fclLaunchKernel','clReleaseEvent') 
-    end if
-
     errcode = clEnqueueNDRangeKernel(cmdq%cl_command_queue, &
                 kernel%cl_kernel, kernel%work_dim, &
                 c_loc(kernel%global_work_offset), &
                 c_loc(kernel%global_work_size), localSizePtr, &
                 cmdq%nDependency, cmdq%dependencyListPtr, &
-                c_loc(cmdQ%lastKernelEvent%cl_event))
+                c_loc(kernelEvent%cl_event))
 
-    call fclDbgWait(cmdQ%lastKernelEvent)
+    call fclDbgWait(kernelEvent)
     call fclPopDependencies(cmdq)
     call fclErrorHandler(errcode,'fclLaunchKernel','clEnqueueNDRangeKernel')
 
-    fclLastKernelEvent = cmdQ%lastKernelEvent
+    fclLastKernelEvent = kernelEvent
+    cmdQ%lastKernelEvent = kernelEvent
 
-    call kernel%pushProfileEvent(cmdQ%lastKernelEvent)
+    call kernel%pushProfileEvent(kernelEvent)
 
   end procedure fclLaunchKernel
   ! ---------------------------------------------------------------------------
@@ -980,21 +976,17 @@
   module procedure fclBarrier_1 !(cmdq)
     !! Enqueue barrier on all events in command queue
     integer(c_int32_t) :: errcode
-
-    ! Decrement event reference counter
-    if (cmdq%lastBarrierEvent%cl_event > 0) then
-      errcode = clReleaseEvent(cmdq%lastBarrierEvent%cl_event)
-      call fclErrorHandler(errcode,'fclBarrier','clReleaseEvent') 
-    end if
+    type(fclEvent), target :: barrierEvent
 
     errcode = clEnqueueBarrierWithWaitList( cmdq%cl_command_queue, &
                   cmdq%nDependency, cmdq%dependencyListPtr , &
-                  c_loc(cmdq%lastBarrierEvent%cl_event))
+                  c_loc(barrierEvent%cl_event))
 
     call fclPopDependencies(cmdq)
     call fclErrorHandler(errcode,'fclBarrierAll','clEnqueueBarrierWithWaitList')
 
-    fclLastBarrierEvent = cmdq%lastBarrierEvent
+    fclLastBarrierEvent = barrierEvent
+    cmdq%lastBarrierEvent = barrierEvent
 
   end procedure fclBarrier_1
   ! ---------------------------------------------------------------------------
@@ -1070,6 +1062,55 @@
   ! ---------------------------------------------------------------------------
 
 
+  module procedure fclEventCopy !(target, source)
+    !! Overloaded assignment for event assignment.
+    !!  Handles opencl reference counting for the underlying event object
+
+    if (target%cl_event > 0) then
+
+      call fclReleaseEvent(target)
+
+    end if
+
+    call fclRetainEvent(source)
+
+    target%cl_event = source%cl_event
+
+  end procedure fclEventCopy
+  ! ---------------------------------------------------------------------------
+
+
+  module procedure fclReleaseEvent !(event)
+    !! Light weight wrapper for clReleaseEvent (decrement reference count)
+    integer(c_int32_t) :: errcode
+
+    if (event%cl_event > 0) then
+
+      errcode = clReleaseEvent(event%cl_event)
+      call fclErrorHandler(errcode,'fclReleaseEvent','clReleaseEvent')
+
+    end if
+
+  end procedure fclReleaseEvent
+  ! ---------------------------------------------------------------------------
+
+
+  module procedure fclRetainEvent !(event)
+    !! Light weight wrapper for clRetainEvent (increment reference count)
+    integer(c_int32_t) :: errcode
+
+
+    if (event%cl_event > 0) then
+
+      errcode = clRetainEvent(event%cl_event)
+      call fclErrorHandler(errcode,'fclRetainEvent','clRetainEvent')
+
+    end if
+
+  end procedure fclRetainEvent
+  ! ---------------------------------------------------------------------------
+
+
   module procedure fclSetDependencyEvent_1 !(cmdq,event,hold)
     !! Specify a single event dependency on specific cmdq
 
@@ -1085,7 +1126,7 @@
     cmdq%nDependency = 1
     cmdq%dependencyListPtr = c_loc(cmdq%dependencyList)
 
-    ! Increment event reference counter
+    ! Explicitly increment event reference counter
     errcode = clRetainEvent(event%cl_event)
     call fclErrorHandler(errcode,'fclSetDependencyEvent','clRetainEvent')
 
@@ -1129,7 +1170,7 @@
     cmdq%nDependency = nEvent
     cmdq%dependencyListPtr = c_loc(cmdq%dependencyList)
 
-    ! Increment event reference counters
+    ! Explicitly increment event reference counters
     do i=1,nEvent
       errcode = clRetainEvent(eventList(i)%cl_event)
       call fclErrorHandler(errcode,'fclSetDependencyEvent','clRetainEvent') 
@@ -1170,7 +1211,7 @@
     integer :: i
     integer(c_int32_t) :: errcode
 
-    ! Decrement event reference counters
+    ! Explicitly decrement event reference counters
     do i=1,cmdq%nDependency
       errcode = clReleaseEvent(cmdq%dependencyList(i))
       call fclErrorHandler(errcode,'fclClearDependencies','clReleaseEvent')