ocl: rework events handling with clSetEventCallback

alalek · alalek · commit daee9821068c · 2017-07-06T13:25:32.000+03:00
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
@@ -1962,7 +1962,7 @@ KernelArg KernelArg::Constant(const Mat& m)
 struct Kernel::Impl
 {
     Impl(const char* kname, const Program& prog) :
-        refcount(1), e(0), nu(0)
+        refcount(1), isInProgress(false), nu(0)
     {
         cl_program ph = (cl_program)prog.ptr();
         cl_int retval = 0;
@@ -2005,11 +2005,15 @@ struct Kernel::Impl
         images.push_back(image);
     }
 
-    void finit()
+    void finit(cl_event e)
     {
+        CV_UNUSED(e);
+#if 0
+        printf("event::callback(%p)\n", e); fflush(stdout);
+#endif
         cleanupUMats();
         images.clear();
-        if(e) { clReleaseEvent(e); e = 0; }
+        isInProgress = false;
         release();
     }
 
@@ -2025,9 +2029,9 @@ struct Kernel::Impl
     cv::String name;
 #endif
     cl_kernel handle;
-    cl_event e;
     enum { MAX_ARRS = 16 };
     UMatData* u[MAX_ARRS];
+    bool isInProgress;
     int nu;
     std::list<Image2D> images;
     bool haveTempDstUMats;
@@ -2037,9 +2041,9 @@ struct Kernel::Impl
 
 extern "C" {
 
-static void CL_CALLBACK oclCleanupCallback(cl_event, cl_int, void *p)
+static void CL_CALLBACK oclCleanupCallback(cl_event e, cl_int, void *p)
 {
-    ((cv::ocl::Kernel::Impl*)p)->finit();
+    ((cv::ocl::Kernel::Impl*)p)->finit(e);
 }
 
 }
@@ -2246,7 +2250,7 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
 {
     CV_INSTRUMENT_REGION_OPENCL_RUN(p->name.c_str());
 
-    if(!p || !p->handle || p->e != 0)
+    if(!p || !p->handle || p->isInProgress)
         return false;
 
     cl_command_queue qq = getQueue(q);
@@ -2265,9 +2269,10 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
         return true;
     if( p->haveTempDstUMats )
         sync = true;
+    cl_event asyncEvent = 0;
     cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims,
                                            offset, globalsize, _localsize, 0, 0,
-                                           sync ? 0 : &p->e);
+                                           sync ? 0 : &asyncEvent);
 #if CV_OPENCL_SHOW_RUN_ERRORS
     if (retval != CL_SUCCESS)
     {
@@ -2283,18 +2288,22 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
     else
     {
         p->addref();
-        CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);
+        p->isInProgress = true;
+        CV_OclDbgAssert(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);
     }
+    if (asyncEvent)
+        clReleaseEvent(asyncEvent);
     return retval == CL_SUCCESS;
 }
 
 bool Kernel::runTask(bool sync, const Queue& q)
 {
-    if(!p || !p->handle || p->e != 0)
+    if(!p || !p->handle || p->isInProgress)
         return false;
 
     cl_command_queue qq = getQueue(q);
-    cl_int retval = clEnqueueTask(qq, p->handle, 0, 0, sync ? 0 : &p->e);
+    cl_event asyncEvent = 0;
+    cl_int retval = clEnqueueTask(qq, p->handle, 0, 0, sync ? 0 : &asyncEvent);
     if( sync || retval != CL_SUCCESS )
     {
         CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS);
@@ -2303,8 +2312,11 @@ bool Kernel::runTask(bool sync, const Queue& q)
     else
     {
         p->addref();
-        CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);
+        p->isInProgress = true;
+        CV_OclDbgAssert(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);
     }
+    if (asyncEvent)
+        clReleaseEvent(asyncEvent);
     return retval == CL_SUCCESS;
 }
 

Original file line number	Diff line number	Diff line change
`@@ -1962,7 +1962,7 @@ KernelArg KernelArg::Constant(const Mat& m)`
`1962`	`1962`	`struct Kernel::Impl`
`1963`	`1963`	`{`
`1964`	`1964`	`Impl(const char* kname, const Program& prog) :`
`1965`		`- refcount(1), e(0), nu(0)`
	`1965`	`+ refcount(1), isInProgress(false), nu(0)`
`1966`	`1966`	`{`
`1967`	`1967`	`cl_program ph = (cl_program)prog.ptr();`
`1968`	`1968`	`cl_int retval = 0;`
`@@ -2005,11 +2005,15 @@ struct Kernel::Impl`
`2005`	`2005`	`images.push_back(image);`
`2006`	`2006`	`}`
`2007`	`2007`
`2008`		`- void finit()`
	`2008`	`+ void finit(cl_event e)`
`2009`	`2009`	`{`
	`2010`	`+ CV_UNUSED(e);`
	`2011`	`+#if 0`
	`2012`	`+ printf("event::callback(%p)\n", e); fflush(stdout);`
	`2013`	`+#endif`
`2010`	`2014`	`cleanupUMats();`
`2011`	`2015`	`images.clear();`
`2012`		`- if(e) { clReleaseEvent(e); e = 0; }`
	`2016`	`+ isInProgress = false;`
`2013`	`2017`	`release();`
`2014`	`2018`	`}`
`2015`	`2019`
`@@ -2025,9 +2029,9 @@ struct Kernel::Impl`
`2025`	`2029`	`cv::String name;`
`2026`	`2030`	`#endif`
`2027`	`2031`	`cl_kernel handle;`
`2028`		`- cl_event e;`
`2029`	`2032`	`enum { MAX_ARRS = 16 };`
`2030`	`2033`	`UMatData* u[MAX_ARRS];`
	`2034`	`+ bool isInProgress;`
`2031`	`2035`	`int nu;`
`2032`	`2036`	`std::list<Image2D> images;`
`2033`	`2037`	`bool haveTempDstUMats;`
`@@ -2037,9 +2041,9 @@ struct Kernel::Impl`
`2037`	`2041`
`2038`	`2042`	`extern "C" {`
`2039`	`2043`
`2040`		`-static void CL_CALLBACK oclCleanupCallback(cl_event, cl_int, void *p)`
	`2044`	`+static void CL_CALLBACK oclCleanupCallback(cl_event e, cl_int, void *p)`
`2041`	`2045`	`{`
`2042`		`- ((cv::ocl::Kernel::Impl*)p)->finit();`
	`2046`	`+ ((cv::ocl::Kernel::Impl*)p)->finit(e);`
`2043`	`2047`	`}`
`2044`	`2048`
`2045`	`2049`	`}`
`@@ -2246,7 +2250,7 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],`
`2246`	`2250`	`{`
`2247`	`2251`	`CV_INSTRUMENT_REGION_OPENCL_RUN(p->name.c_str());`
`2248`	`2252`
`2249`		`- if(!p \|\| !p->handle \|\| p->e != 0)`
	`2253`	`+ if(!p \|\| !p->handle \|\| p->isInProgress)`
`2250`	`2254`	`return false;`
`2251`	`2255`
`2252`	`2256`	`cl_command_queue qq = getQueue(q);`
`@@ -2265,9 +2269,10 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],`
`2265`	`2269`	`return true;`
`2266`	`2270`	`if( p->haveTempDstUMats )`
`2267`	`2271`	`sync = true;`
	`2272`	`+ cl_event asyncEvent = 0;`
`2268`	`2273`	`cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims,`
`2269`	`2274`	`offset, globalsize, _localsize, 0, 0,`
`2270`		`- sync ? 0 : &p->e);`
	`2275`	`+ sync ? 0 : &asyncEvent);`
`2271`	`2276`	`#if CV_OPENCL_SHOW_RUN_ERRORS`
`2272`	`2277`	`if (retval != CL_SUCCESS)`
`2273`	`2278`	`{`
`@@ -2283,18 +2288,22 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],`
`2283`	`2288`	`else`
`2284`	`2289`	`{`
`2285`	`2290`	`p->addref();`
`2286`		`- CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);`
	`2291`	`+ p->isInProgress = true;`
	`2292`	`+ CV_OclDbgAssert(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);`
`2287`	`2293`	`}`
	`2294`	`+ if (asyncEvent)`
	`2295`	`+ clReleaseEvent(asyncEvent);`
`2288`	`2296`	`return retval == CL_SUCCESS;`
`2289`	`2297`	`}`
`2290`	`2298`
`2291`	`2299`	`bool Kernel::runTask(bool sync, const Queue& q)`
`2292`	`2300`	`{`
`2293`		`- if(!p \|\| !p->handle \|\| p->e != 0)`
	`2301`	`+ if(!p \|\| !p->handle \|\| p->isInProgress)`
`2294`	`2302`	`return false;`
`2295`	`2303`
`2296`	`2304`	`cl_command_queue qq = getQueue(q);`
`2297`		`- cl_int retval = clEnqueueTask(qq, p->handle, 0, 0, sync ? 0 : &p->e);`
	`2305`	`+ cl_event asyncEvent = 0;`
	`2306`	`+ cl_int retval = clEnqueueTask(qq, p->handle, 0, 0, sync ? 0 : &asyncEvent);`
`2298`	`2307`	`if( sync \|\| retval != CL_SUCCESS )`
`2299`	`2308`	`{`
`2300`	`2309`	`CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS);`
`@@ -2303,8 +2312,11 @@ bool Kernel::runTask(bool sync, const Queue& q)`
`2303`	`2312`	`else`
`2304`	`2313`	`{`
`2305`	`2314`	`p->addref();`
`2306`		`- CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);`
	`2315`	`+ p->isInProgress = true;`
	`2316`	`+ CV_OclDbgAssert(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);`
`2307`	`2317`	`}`
	`2318`	`+ if (asyncEvent)`
	`2319`	`+ clReleaseEvent(asyncEvent);`
`2308`	`2320`	`return retval == CL_SUCCESS;`
`2309`	`2321`	`}`
`2310`	`2322`