diff --git a/src/Focal.f90 b/src/Focal.f90 index 801d2d2..1a28750 100644 --- a/src/Focal.f90 +++ b/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. diff --git a/src/Focal_Memory.f90 b/src/Focal_Memory.f90 index 7c641fa..c59e6e2 100644 --- a/src/Focal_Memory.f90 +++ b/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') diff --git a/src/Focal_Profile.f90 b/src/Focal_Profile.f90 index 0309c8f..d52322e 100644 --- a/src/Focal_Profile.f90 +++ b/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 diff --git a/src/Focal_Setup.f90 b/src/Focal_Setup.f90 index 29de831..9433ca3 100644 --- a/src/Focal_Setup.f90 +++ b/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')