Fix an issue with Kernel object reset release when consecutive Kernel::run calls

insoow · insoow · commit cc7f9f54695d · 2017-01-30T15:20:28.000-08:00
Kernel::run launch OCL gpu kernels and set a event callback function
to decreate the ref count of UMat or remove UMat when the lauched workloads
are completed. However, for some OCL kernels requires multiple call of
Kernel::run function with some kernel parameter changes (e.g., input
and output buffer offset) to get the final computation result.
In the case, the current implementation requires unnecessary
synchronization and cleanupMat.

This fix requires the user to specify whether there will be more work or not.
If there is no remaining computation, the Kernel::run will reset the
kernel object

Signed-off-by: Woo, Insoo &lt;insoo.woo@intel.com&gt;
diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp
@@ -578,9 +578,17 @@ class CV_EXPORTS Kernel
     @param localsize work-group size for each dimension.
     @param sync specify whether to wait for OpenCL computation to finish before return.
     @param q command queue
+    @param moreWorkDone specify whether there will the remaining work to be computed (more Kernel::run calls).
+    When a computation requires multiple kernel execution by changing input and output buffer offset to get
+    the final computation results.
+    kernel.setArg(0, ..);
+    kernel.setArg(1, offset);
+    kernel.run(..., q, true);
+    kernel.setArg(1, offset+256);
+    kernel.run(..., q, false);
     */
     bool run(int dims, size_t globalsize[],
-             size_t localsize[], bool sync, const Queue& q=Queue());
+             size_t localsize[], bool sync, const Queue& q=Queue(), bool moreWorkDone = false);
     bool runTask(bool sync, const Queue& q=Queue());
 
     size_t workGroupSize() const;
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
@@ -3185,7 +3185,7 @@ struct Kernel::Impl
 
     void cleanupUMats()
     {
-        for( int i = 0; i < MAX_ARRS; i++ )
+        for( int i = 0; i < nu; i++ )
             if( u[i] )
             {
                 if( CV_XADD(&u[i]->urefcount, -1) == 1 )
@@ -3446,9 +3446,8 @@ int Kernel::set(int i, const KernelArg& arg)
     return i+1;
 }
 
-
 bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
-                 bool sync, const Queue& q)
+                 bool sync, const Queue& q, bool moreWorkDone)
 {
     CV_INSTRUMENT_REGION_OPENCL_RUN(p->name.c_str());
 
@@ -3469,11 +3468,11 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
     }
     if( total == 0 )
         return true;
-    if( p->haveTempDstUMats )
+    if( p->haveTempDstUMats && !moreWorkDone)
         sync = true;
     cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims,
                                            offset, globalsize, _localsize, 0, 0,
-                                           sync ? 0 : &p->e);
+                                           sync ? 0 : (moreWorkDone? 0: &p->e ));
 #if CV_OPENCL_SHOW_RUN_ERRORS
     if (retval != CL_SUCCESS)
     {
@@ -3484,9 +3483,9 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
     if( sync || retval != CL_SUCCESS )
     {
         CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS);
-        p->cleanupUMats();
+        if (!moreWorkDone) p->cleanupUMats();
     }
-    else
+    else if (!moreWorkDone)
     {
         p->addref();
         CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);

Original file line number	Diff line number	Diff line change
`@@ -3185,7 +3185,7 @@ struct Kernel::Impl`
`3185`	`3185`
`3186`	`3186`	`void cleanupUMats()`
`3187`	`3187`	`{`
`3188`		`- for( int i = 0; i < MAX_ARRS; i++ )`
	`3188`	`+ for( int i = 0; i < nu; i++ )`
`3189`	`3189`	`if( u[i] )`
`3190`	`3190`	`{`
`3191`	`3191`	`if( CV_XADD(&u[i]->urefcount, -1) == 1 )`
`@@ -3446,9 +3446,8 @@ int Kernel::set(int i, const KernelArg& arg)`
`3446`	`3446`	`return i+1;`
`3447`	`3447`	`}`
`3448`	`3448`
`3449`		`-`
`3450`	`3449`	`bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],`
`3451`		`- bool sync, const Queue& q)`
	`3450`	`+ bool sync, const Queue& q, bool moreWorkDone)`
`3452`	`3451`	`{`
`3453`	`3452`	`CV_INSTRUMENT_REGION_OPENCL_RUN(p->name.c_str());`
`3454`	`3453`
`@@ -3469,11 +3468,11 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],`
`3469`	`3468`	`}`
`3470`	`3469`	`if( total == 0 )`
`3471`	`3470`	`return true;`
`3472`		`- if( p->haveTempDstUMats )`
	`3471`	`+ if( p->haveTempDstUMats && !moreWorkDone)`
`3473`	`3472`	`sync = true;`
`3474`	`3473`	`cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims,`
`3475`	`3474`	`offset, globalsize, _localsize, 0, 0,`
`3476`		`- sync ? 0 : &p->e);`
	`3475`	`+ sync ? 0 : (moreWorkDone? 0: &p->e ));`
`3477`	`3476`	`#if CV_OPENCL_SHOW_RUN_ERRORS`
`3478`	`3477`	`if (retval != CL_SUCCESS)`
`3479`	`3478`	`{`
`@@ -3484,9 +3483,9 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],`
`3484`	`3483`	`if( sync \|\| retval != CL_SUCCESS )`
`3485`	`3484`	`{`
`3486`	`3485`	`CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS);`
`3487`		`- p->cleanupUMats();`
	`3486`	`+ if (!moreWorkDone) p->cleanupUMats();`
`3488`	`3487`	`}`
`3489`		`- else`
	`3488`	`+ else if (!moreWorkDone)`
`3490`	`3489`	`{`
`3491`	`3490`	`p->addref();`
`3492`	`3491`	`CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);`