Merge pull request #9988 from alalek:ocl_verbose_api_errors

modules/core/src/ocl.cpp

@@ -57,7 +57,10 @@
 #include "opencl_kernels_core.hpp"
 
 #define CV_OPENCL_ALWAYS_SHOW_BUILD_LOG 0
-#define CV_OPENCL_SHOW_RUN_ERRORS       0
+
+#define CV_OPENCL_SHOW_RUN_KERNELS      0
+#define CV_OPENCL_TRACE_CHECK           0
+
 #define CV_OPENCL_SHOW_SVM_ERROR_LOG    1
 #define CV_OPENCL_SHOW_SVM_LOG          0
 
@@ -94,9 +97,15 @@
 #include "ocl_deprecated.hpp"
 #endif // HAVE_OPENCL
 
-#ifdef _DEBUG
-#define CV_OclDbgAssert CV_DbgAssert
-#else
+#ifdef HAVE_OPENCL_SVM
+#include "opencv2/core/opencl/runtime/opencl_svm_20.hpp"
+#include "opencv2/core/opencl/runtime/opencl_svm_hsa_extension.hpp"
+#include "opencv2/core/opencl/opencl_svm.hpp"
+#endif
+
+namespace cv { namespace ocl {
+
+#ifndef _DEBUG
 static bool isRaiseError()
 {
     static bool initialized = false;
@@ -108,16 +117,55 @@ static bool isRaiseError()
     }
     return value;
 }
-#define CV_OclDbgAssert(expr) do { if (isRaiseError()) { CV_Assert(expr); } else { (void)(expr); } } while ((void)0, 0)
 #endif
 
-#ifdef HAVE_OPENCL_SVM
-#include "opencv2/core/opencl/runtime/opencl_svm_20.hpp"
-#include "opencv2/core/opencl/runtime/opencl_svm_hsa_extension.hpp"
-#include "opencv2/core/opencl/opencl_svm.hpp"
+#if CV_OPENCL_TRACE_CHECK
+static inline
+void traceOpenCLCheck(cl_int status, const char* message)
+{
+    std::cout << "OpenCV(OpenCL:" << status << "): " << message << std::endl << std::flush;
+}
+#define CV_OCL_TRACE_CHECK_RESULT(status, message) traceOpenCLCheck(status, message)
+#else
+#define CV_OCL_TRACE_CHECK_RESULT(status, message) /* nothing */
 #endif
 
-namespace cv { namespace ocl {
+#define CV_OCL_API_ERROR_MSG(check_result, msg) \
+    cv::format("OpenCL error %s (%d) during call: %s", getOpenCLErrorString(check_result), check_result, msg)
+
+#define CV_OCL_CHECK_RESULT(check_result, msg) \
+    do { \
+        CV_OCL_TRACE_CHECK_RESULT(check_result, msg); \
+        if (check_result != CL_SUCCESS) \
+        { \
+            if (0) { const char* msg_ = (msg); (void)msg_; /* ensure const char* type (cv::String without c_str()) */ } \
+            cv::String error_msg = CV_OCL_API_ERROR_MSG(check_result, msg); \
+            CV_Error(Error::OpenCLApiCallError, error_msg); \
+        } \
+    } while (0)
+
+#define CV_OCL_CHECK_(expr, check_result) do { expr; CV_OCL_CHECK_RESULT(check_result, #expr); } while (0)
+
+#define CV_OCL_CHECK(expr) do { cl_int __cl_result = (expr); CV_OCL_CHECK_RESULT(__cl_result, #expr); } while (0)
+
+#ifdef _DEBUG
+#define CV_OCL_DBG_CHECK_RESULT(check_result, msg) CV_OCL_CHECK_RESULT(check_result, msg)
+#define CV_OCL_DBG_CHECK(expr) CV_OCL_CHECK(expr)
+#define CV_OCL_DBG_CHECK_(expr, check_result) CV_OCL_CHECK_(expr, check_result)
+#else
+#define CV_OCL_DBG_CHECK_RESULT(check_result, msg) \
+    do { \
+        CV_OCL_TRACE_CHECK_RESULT(check_result, msg); \
+        if (check_result != CL_SUCCESS && isRaiseError()) \
+        { \
+            if (0) { const char* msg_ = (msg); (void)msg_; /* ensure const char* type (cv::String without c_str()) */ } \
+            cv::String error_msg = CV_OCL_API_ERROR_MSG(check_result, msg); \
+            CV_Error(Error::OpenCLApiCallError, error_msg); \
+        } \
+    } while (0)
+#define CV_OCL_DBG_CHECK_(expr, check_result) do { expr; CV_OCL_DBG_CHECK_RESULT(check_result, #expr); } while (0)
+#define CV_OCL_DBG_CHECK(expr) do { cl_int __cl_result = (expr); CV_OCL_DBG_CHECK_RESULT(__cl_result, #expr); } while (0)
+#endif
 
 struct UMat2D
 {
@@ -428,7 +476,7 @@ struct Platform::Impl
             {
                 char buf[1000];
                 size_t len = 0;
-                CV_OclDbgAssert(clGetPlatformInfo(handle, CL_PLATFORM_VENDOR, sizeof(buf), buf, &len) == CL_SUCCESS);
+                CV_OCL_DBG_CHECK(clGetPlatformInfo(handle, CL_PLATFORM_VENDOR, sizeof(buf), buf, &len));
                 buf[len] = '\0';
                 vendor = String(buf);
             }
@@ -856,8 +904,8 @@ void Device::maxWorkItemSizes(size_t* sizes) const
     {
         const int MAX_DIMS = 32;
         size_t retsz = 0;
-        CV_OclDbgAssert(clGetDeviceInfo(p->handle, CL_DEVICE_MAX_WORK_ITEM_SIZES,
-                MAX_DIMS*sizeof(sizes[0]), &sizes[0], &retsz) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clGetDeviceInfo(p->handle, CL_DEVICE_MAX_WORK_ITEM_SIZES,
+                MAX_DIMS*sizeof(sizes[0]), &sizes[0], &retsz));
     }
 }
 
@@ -1042,12 +1090,12 @@ static cl_device_id selectOpenCLDevice()
     std::vector<cl_platform_id> platforms;
     {
         cl_uint numPlatforms = 0;
-        CV_OclDbgAssert(clGetPlatformIDs(0, NULL, &numPlatforms) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clGetPlatformIDs(0, NULL, &numPlatforms));
 
         if (numPlatforms == 0)
             return NULL;
         platforms.resize((size_t)numPlatforms);
-        CV_OclDbgAssert(clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms));
         platforms.resize(numPlatforms);
     }
 
@@ -1057,7 +1105,7 @@ static cl_device_id selectOpenCLDevice()
         for (size_t i = 0; i < platforms.size(); i++)
         {
             std::string name;
-            CV_OclDbgAssert(getStringInfo(clGetPlatformInfo, platforms[i], CL_PLATFORM_NAME, name) == CL_SUCCESS);
+            CV_OCL_DBG_CHECK(getStringInfo(clGetPlatformInfo, platforms[i], CL_PLATFORM_NAME, name));
             if (name.find(platform) != std::string::npos)
             {
                 selectedPlatform = (int)i;
@@ -1108,13 +1156,19 @@ static cl_device_id selectOpenCLDevice()
         {
             cl_uint count = 0;
             cl_int status = clGetDeviceIDs(platforms[i], deviceType, 0, NULL, &count);
-            CV_OclDbgAssert(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND);
+            if (!(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND))
+            {
+                CV_OCL_DBG_CHECK_RESULT(status, "clGetDeviceIDs get count");
+            }
             if (count == 0)
                 continue;
             size_t base = devices.size();
             devices.resize(base + count);
             status = clGetDeviceIDs(platforms[i], deviceType, count, &devices[base], &count);
-            CV_OclDbgAssert(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND);
+            if (!(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND))
+            {
+                CV_OCL_DBG_CHECK_RESULT(status, "clGetDeviceIDs get IDs");
+            }
         }
 
         for (size_t i = (isID ? deviceID : 0);
@@ -1122,12 +1176,12 @@ static cl_device_id selectOpenCLDevice()
              i++)
         {
             std::string name;
-            CV_OclDbgAssert(getStringInfo(clGetDeviceInfo, devices[i], CL_DEVICE_NAME, name) == CL_SUCCESS);
+            CV_OCL_DBG_CHECK(getStringInfo(clGetDeviceInfo, devices[i], CL_DEVICE_NAME, name));
             cl_bool useGPU = true;
             if(tempStrDeviceType == "dgpu" || tempStrDeviceType == "igpu")
             {
                 cl_bool isIGPU = CL_FALSE;
-                clGetDeviceInfo(devices[i], CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(isIGPU), &isIGPU, NULL);
+                CV_OCL_DBG_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(isIGPU), &isIGPU, NULL));
                 useGPU = tempStrDeviceType == "dgpu" ? !isIGPU : isIGPU;
             }
             if ( (isID || name.find(deviceName) != std::string::npos) && useGPU)
@@ -1257,7 +1311,7 @@ struct Context::Impl
             return;
 
         cl_platform_id pl = NULL;
-        CV_OclDbgAssert(clGetDeviceInfo(d, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &pl, NULL) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clGetDeviceInfo(d, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &pl, NULL));
 
         cl_context_properties prop[] =
         {
@@ -1270,6 +1324,7 @@ struct Context::Impl
         cl_int status;
 
         handle = clCreateContext(prop, nd, &d, 0, 0, &status);
+        CV_OCL_DBG_CHECK_RESULT(status, "clCreateContext");
 
         bool ok = handle != 0 && status == CL_SUCCESS;
         if( ok )
@@ -1295,12 +1350,12 @@ struct Context::Impl
 
         cl_uint i, nd0 = 0, nd = 0;
         int dtype = dtype0 & 15;
-        CV_OclDbgAssert(clGetDeviceIDs( pl, dtype, 0, 0, &nd0 ) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clGetDeviceIDs(pl, dtype, 0, 0, &nd0));
 
         AutoBuffer<void*> dlistbuf(nd0*2+1);
         cl_device_id* dlist = (cl_device_id*)(void**)dlistbuf;
         cl_device_id* dlist_new = dlist + nd0;
-        CV_OclDbgAssert(clGetDeviceIDs( pl, dtype, nd0, dlist, &nd0 ) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clGetDeviceIDs(pl, dtype, nd0, dlist, &nd0));
         String name0;
 
         for(i = 0; i < nd0; i++)
@@ -1326,6 +1381,7 @@ struct Context::Impl
         nd = 1;
 
         handle = clCreateContext(prop, nd, dlist_new, 0, 0, &retval);
+        CV_OCL_DBG_CHECK_RESULT(retval, "clCreateContext");
         bool ok = handle != 0 && retval == CL_SUCCESS;
         if( ok )
         {
@@ -1339,7 +1395,7 @@ struct Context::Impl
     {
         if(handle)
         {
-            clReleaseContext(handle);
+            CV_OCL_DBG_CHECK(clReleaseContext(handle));
             handle = NULL;
         }
         devices.clear();
@@ -1527,8 +1583,7 @@ struct Context::Impl
                     goto noSVM;
                 }
                 cl_platform_id p = NULL;
-                status = clGetDeviceInfo((cl_device_id)device.ptr(), CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &p, NULL);
-                CV_Assert(status == CL_SUCCESS);
+                CV_OCL_CHECK(status = clGetDeviceInfo((cl_device_id)device.ptr(), CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &p, NULL));
                 svmFunctions.fn_clSVMAlloc = (clSVMAllocAMD_fn)clGetExtensionFunctionAddressForPlatform(p, "clSVMAllocAMD");
                 svmFunctions.fn_clSVMFree = (clSVMFreeAMD_fn)clGetExtensionFunctionAddressForPlatform(p, "clSVMFreeAMD");
                 svmFunctions.fn_clSetKernelArgSVMPointer = (clSetKernelArgSVMPointerAMD_fn)clGetExtensionFunctionAddressForPlatform(p, "clSetKernelArgSVMPointerAMD");
@@ -1748,13 +1803,11 @@ static void get_platform_name(cl_platform_id id, String& name)
 {
     // get platform name string length
     size_t sz = 0;
-    if (CL_SUCCESS != clGetPlatformInfo(id, CL_PLATFORM_NAME, 0, 0, &sz))
-        CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "clGetPlatformInfo failed!");
+    CV_OCL_CHECK(clGetPlatformInfo(id, CL_PLATFORM_NAME, 0, 0, &sz));
 
     // get platform name string
     AutoBuffer<char> buf(sz + 1);
-    if (CL_SUCCESS != clGetPlatformInfo(id, CL_PLATFORM_NAME, sz, buf, 0))
-        CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "clGetPlatformInfo failed!");
+    CV_OCL_CHECK(clGetPlatformInfo(id, CL_PLATFORM_NAME, sz, buf, 0));
 
     // just in case, ensure trailing zero for ASCIIZ string
     buf[sz] = 0;
@@ -1769,16 +1822,14 @@ void attachContext(const String& platformName, void* platformID, void* context,
 {
     cl_uint cnt = 0;
 
-    if(CL_SUCCESS != clGetPlatformIDs(0, 0, &cnt))
-        CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "clGetPlatformIDs failed!");
+    CV_OCL_CHECK(clGetPlatformIDs(0, 0, &cnt));
 
     if (cnt == 0)
         CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "no OpenCL platform available!");
 
     std::vector<cl_platform_id> platforms(cnt);
 
-    if(CL_SUCCESS != clGetPlatformIDs(cnt, &platforms[0], 0))
-        CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "clGetPlatformIDs failed!");
+    CV_OCL_CHECK(clGetPlatformIDs(cnt, &platforms[0], 0));
 
     bool platformAvailable = false;
 
@@ -1810,8 +1861,7 @@ void attachContext(const String& platformName, void* platformID, void* context,
     // attach supplied context to OpenCV
     initializeContextFromHandle(ctx, platformID, context, deviceID);
 
-    if(CL_SUCCESS != clRetainContext((cl_context)context))
-        CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "clRetainContext failed!");
+    CV_OCL_CHECK(clRetainContext((cl_context)context));
 
     // clear command queue, if any
     getCoreTlsData().get()->oclQueue.finish();
@@ -1831,7 +1881,7 @@ void initializeContextFromHandle(Context& ctx, void* platform, void* _context, v
     Context::Impl * impl = ctx.p;
     if (impl->handle)
     {
-        CV_OclDbgAssert(clReleaseContext(impl->handle) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clReleaseContext(impl->handle));
     }
     impl->devices.clear();
 
@@ -1861,8 +1911,7 @@ struct Queue::Impl
         handle = q;
 
         cl_command_queue_properties props = 0;
-        cl_int result = clGetCommandQueueInfo(handle, CL_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &props, NULL);
-        CV_Assert(result && "clGetCommandQueueInfo(CL_QUEUE_PROPERTIES)");
+        CV_OCL_CHECK(clGetCommandQueueInfo(handle, CL_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &props, NULL));
         isProfilingQueue_ = !!(props & CL_QUEUE_PROFILING_ENABLE);
     }
 
@@ -1889,8 +1938,7 @@ struct Queue::Impl
             dh = (cl_device_id)pc->device(0).ptr();
         cl_int retval = 0;
         cl_command_queue_properties props = withProfiling ? CL_QUEUE_PROFILING_ENABLE : 0;
-        handle = clCreateCommandQueue(ch, dh, props, &retval);
-        CV_OclDbgAssert(retval == CL_SUCCESS);
+        CV_OCL_DBG_CHECK_(handle = clCreateCommandQueue(ch, dh, props, &retval), retval);
         isProfilingQueue_ = withProfiling;
     }
 
@@ -1902,8 +1950,8 @@ struct Queue::Impl
         {
             if(handle)
             {
-                clFinish(handle);
-                clReleaseCommandQueue(handle);
+                CV_OCL_DBG_CHECK(clFinish(handle));
+                CV_OCL_DBG_CHECK(clReleaseCommandQueue(handle));
                 handle = NULL;
             }
         }
@@ -1918,15 +1966,15 @@ struct Queue::Impl
             return profiling_queue_;
 
         cl_context ctx = 0;
-        CV_Assert(CL_SUCCESS == clGetCommandQueueInfo(handle, CL_QUEUE_CONTEXT, sizeof(cl_context), &ctx, NULL));
+        CV_OCL_CHECK(clGetCommandQueueInfo(handle, CL_QUEUE_CONTEXT, sizeof(cl_context), &ctx, NULL));
 
         cl_device_id device = 0;
-        CV_Assert(CL_SUCCESS == clGetCommandQueueInfo(handle, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device, NULL));
+        CV_OCL_CHECK(clGetCommandQueueInfo(handle, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device, NULL));
 
         cl_int result = CL_SUCCESS;
         cl_command_queue_properties props = CL_QUEUE_PROFILING_ENABLE;
         cl_command_queue q = clCreateCommandQueue(ctx, device, props, &result);
-        CV_Assert(result == CL_SUCCESS && "clCreateCommandQueue(with CL_QUEUE_PROFILING_ENABLE)");
+        CV_OCL_DBG_CHECK_RESULT(result, "clCreateCommandQueue(with CL_QUEUE_PROFILING_ENABLE)");
 
         Queue queue;
         queue.p = new Impl(q, true);
@@ -1989,7 +2037,7 @@ void Queue::finish()
 {
     if(p && p->handle)
     {
-        CV_OclDbgAssert(clFinish(p->handle) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clFinish(p->handle));
     }
 }
 
@@ -2044,16 +2092,16 @@ KernelArg KernelArg::Constant(const Mat& m)
 struct Kernel::Impl
 {
     Impl(const char* kname, const Program& prog) :
-        refcount(1), isInProgress(false), nu(0)
+        refcount(1), handle(NULL), isInProgress(false), nu(0)
     {
         cl_program ph = (cl_program)prog.ptr();
         cl_int retval = 0;
-#ifdef ENABLE_INSTRUMENTATION
         name = kname;
-#endif
-        handle = ph != 0 ?
-            clCreateKernel(ph, kname, &retval) : 0;
-        CV_OclDbgAssert(retval == CL_SUCCESS);
+        if (ph)
+        {
+            handle = clCreateKernel(ph, kname, &retval);
+            CV_OCL_DBG_CHECK_RESULT(retval, cv::format("clCreateKernel('%s')", kname).c_str());
+        }
         for( int i = 0; i < MAX_ARRS; i++ )
             u[i] = 0;
         haveTempDstUMats = false;
@@ -2093,9 +2141,6 @@ struct Kernel::Impl
     void finit(cl_event e)
     {
         CV_UNUSED(e);
-#if 0
-        printf("event::callback(%p)\n", e); fflush(stdout);
-#endif
         cleanupUMats();
         images.clear();
         isInProgress = false;
@@ -2108,14 +2153,14 @@ struct Kernel::Impl
     ~Impl()
     {
         if(handle)
-            clReleaseKernel(handle);
+        {
+            CV_OCL_DBG_CHECK(clReleaseKernel(handle));
+        }
     }
 
     IMPLEMENT_REFCOUNTABLE();
 
-#ifdef ENABLE_INSTRUMENTATION
     cv::String name;
-#endif
     cl_kernel handle;
     enum { MAX_ARRS = 16 };
     UMatData* u[MAX_ARRS];
@@ -2230,7 +2275,7 @@ int Kernel::set(int i, const void* value, size_t sz)
         p->cleanupUMats();
 
     cl_int retval = clSetKernelArg(p->handle, (cl_uint)i, sz, value);
-    CV_OclDbgAssert(retval == CL_SUCCESS);
+    CV_OCL_DBG_CHECK_RESULT(retval, cv::format("clSetKernelArg('%s', arg_index=%d, size=%d, value=%p)", p->name.c_str(), (int)i, (int)sz, (void*)value).c_str());
     if (retval != CL_SUCCESS)
         return -1;
     return i+1;
@@ -2256,6 +2301,7 @@ int Kernel::set(int i, const KernelArg& arg)
         return i;
     if( i == 0 )
         p->cleanupUMats();
+    cl_int status = 0;
     if( arg.m )
     {
         int accessFlags = ((arg.flags & KernelArg::READ_ONLY) ? ACCESS_READ : 0) +
@@ -2278,16 +2324,17 @@ int Kernel::set(int i, const KernelArg& arg)
             uchar*& svmDataPtr = (uchar*&)arg.m->u->handle;
             CV_OPENCL_SVM_TRACE_P("clSetKernelArgSVMPointer: %p\n", svmDataPtr);
 #if 1 // TODO
-            cl_int status = svmFns->fn_clSetKernelArgSVMPointer(p->handle, (cl_uint)i, svmDataPtr);
+            status = svmFns->fn_clSetKernelArgSVMPointer(p->handle, (cl_uint)i, svmDataPtr);
 #else
-            cl_int status = svmFns->fn_clSetKernelArgSVMPointer(p->handle, (cl_uint)i, &svmDataPtr);
+            status = svmFns->fn_clSetKernelArgSVMPointer(p->handle, (cl_uint)i, &svmDataPtr);
 #endif
-            CV_Assert(status == CL_SUCCESS);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArgSVMPointer('%s', arg_index=%d, ptr=%p)", p->name.c_str(), (int)i, (void*)svmDataPtr).c_str());
         }
         else
 #endif
         {
-            CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h) == CL_SUCCESS);
+            status = clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, cl_mem=%p)", p->name.c_str(), (int)i, (void*)h).c_str());
         }
 
         if (ptronly)
@@ -2297,38 +2344,49 @@ int Kernel::set(int i, const KernelArg& arg)
         else if( arg.m->dims <= 2 )
         {
             UMat2D u2d(*arg.m);
-            CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u2d.step), &u2d.step) == CL_SUCCESS);
-            CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u2d.offset), &u2d.offset) == CL_SUCCESS);
+            status = clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u2d.step), &u2d.step);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, step_value=%d)", p->name.c_str(), (int)(i+1), (int)u2d.step).c_str());
+            status = clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u2d.offset), &u2d.offset);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, offset_value=%d)", p->name.c_str(), (int)(i+2), (int)u2d.offset).c_str());
             i += 3;
 
             if( !(arg.flags & KernelArg::NO_SIZE) )
             {
                 int cols = u2d.cols*arg.wscale/arg.iwscale;
-                CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, sizeof(u2d.rows), &u2d.rows) == CL_SUCCESS);
-                CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(cols), &cols) == CL_SUCCESS);
+                status = clSetKernelArg(p->handle, (cl_uint)i, sizeof(u2d.rows), &u2d.rows);
+                CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, rows_value=%d)", p->name.c_str(), (int)i, (int)u2d.rows).c_str());
+                status = clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(cols), &cols);
+                CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, cols_value=%d)", p->name.c_str(), (int)(i+1), (int)cols).c_str());
                 i += 2;
             }
         }
         else
         {
             UMat3D u3d(*arg.m);
-            CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u3d.slicestep), &u3d.slicestep) == CL_SUCCESS);
-            CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u3d.step), &u3d.step) == CL_SUCCESS);
-            CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+3), sizeof(u3d.offset), &u3d.offset) == CL_SUCCESS);
+            status = clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u3d.slicestep), &u3d.slicestep);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, slicestep_value=%d)", p->name.c_str(), (int)(i+1), (int)u3d.slicestep).c_str());
+            status = clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u3d.step), &u3d.step);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, step_value=%d)", p->name.c_str(), (int)(i+2), (int)u3d.step).c_str());
+            status = clSetKernelArg(p->handle, (cl_uint)(i+3), sizeof(u3d.offset), &u3d.offset);
+            CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, offset_value=%d)", p->name.c_str(), (int)(i+3), (int)u3d.offset).c_str());
             i += 4;
             if( !(arg.flags & KernelArg::NO_SIZE) )
             {
                 int cols = u3d.cols*arg.wscale/arg.iwscale;
-                CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, sizeof(u3d.slices), &u3d.slices) == CL_SUCCESS);
-                CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u3d.rows), &u3d.rows) == CL_SUCCESS);
-                CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u3d.cols), &cols) == CL_SUCCESS);
+                status = clSetKernelArg(p->handle, (cl_uint)i, sizeof(u3d.slices), &u3d.slices);
+                CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, slices_value=%d)", p->name.c_str(), (int)i, (int)u3d.slices).c_str());
+                status = clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u3d.rows), &u3d.rows);
+                CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, rows_value=%d)", p->name.c_str(), (int)(i+1), (int)u3d.rows).c_str());
+                status = clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u3d.cols), &cols);
+                CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, cols_value=%d)", p->name.c_str(), (int)(i+2), (int)cols).c_str());
                 i += 3;
             }
         }
         p->addUMat(*arg.m, (accessFlags & ACCESS_WRITE) != 0);
         return i;
     }
-    CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, arg.sz, arg.obj) == CL_SUCCESS);
+    status = clSetKernelArg(p->handle, (cl_uint)i, arg.sz, arg.obj);
+    CV_OCL_DBG_CHECK_RESULT(status, cv::format("clSetKernelArg('%s', arg_index=%d, size=%d, obj=%p)", p->name.c_str(), (int)i, (int)arg.sz, (void*)arg.obj).c_str());
     return i+1;
 }
 
@@ -2360,7 +2418,7 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
 bool Kernel::Impl::run(int dims, size_t globalsize[], size_t localsize[],
         bool sync, int64* timeNS, const Queue& q)
 {
-    CV_INSTRUMENT_REGION_OPENCL_RUN(p->name.c_str());
+    CV_INSTRUMENT_REGION_OPENCL_RUN(name.c_str());
 
     if (!handle || isInProgress)
         return false;
@@ -2374,24 +2432,37 @@ bool Kernel::Impl::run(int dims, size_t globalsize[], size_t localsize[],
     cl_int retval = clEnqueueNDRangeKernel(qq, handle, (cl_uint)dims,
                                            NULL, globalsize, localsize, 0, 0,
                                            (sync && !timeNS) ? 0 : &asyncEvent);
-#if CV_OPENCL_SHOW_RUN_ERRORS
+#if !CV_OPENCL_SHOW_RUN_KERNELS
     if (retval != CL_SUCCESS)
+#endif
     {
-        printf("OpenCL program returns error: %d\n", retval);
-        fflush(stdout);
+        cv::String msg = cv::format("clEnqueueNDRangeKernel('%s', dims=%d, globalsize=%dx%dx%d, localsize=%s) sync=%s", name.c_str(), (int)dims,
+                        globalsize[0], (dims > 1 ? globalsize[1] : 1), (dims > 2 ? globalsize[2] : 1),
+                        (localsize ? cv::format("%dx%dx%d", localsize[0], (dims > 1 ? localsize[1] : 1), (dims > 2 ? localsize[2] : 1)) : cv::String("NULL")).c_str(),
+                        sync ? "true" : "false"
+                        );
+        if (retval != CL_SUCCESS)
+        {
+            msg = CV_OCL_API_ERROR_MSG(retval, msg.c_str());
         }
+#if CV_OPENCL_TRACE_CHECK
+        CV_OCL_TRACE_CHECK_RESULT(retval, msg.c_str());
+#else
+        printf("%s\n", msg.c_str());
+        fflush(stdout);
 #endif
+    }
     if (sync || retval != CL_SUCCESS)
     {
-        CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clFinish(qq));
         if (timeNS)
         {
             if (retval == CL_SUCCESS)
             {
-                clWaitForEvents(1, &asyncEvent);
+                CV_OCL_DBG_CHECK(clWaitForEvents(1, &asyncEvent));
                 cl_ulong startTime, stopTime;
-                CV_Assert(CL_SUCCESS == clGetEventProfilingInfo(asyncEvent, CL_PROFILING_COMMAND_START, sizeof(startTime), &startTime, NULL));
-                CV_Assert(CL_SUCCESS == clGetEventProfilingInfo(asyncEvent, CL_PROFILING_COMMAND_END, sizeof(stopTime), &stopTime, NULL));
+                CV_OCL_CHECK(clGetEventProfilingInfo(asyncEvent, CL_PROFILING_COMMAND_START, sizeof(startTime), &startTime, NULL));
+                CV_OCL_CHECK(clGetEventProfilingInfo(asyncEvent, CL_PROFILING_COMMAND_END, sizeof(stopTime), &stopTime, NULL));
                 *timeNS = (int64)(stopTime - startTime);
             }
             else
@@ -2405,10 +2476,10 @@ bool Kernel::Impl::run(int dims, size_t globalsize[], size_t localsize[],
     {
         addref();
         isInProgress = true;
-        CV_OclDbgAssert(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, this) == CL_SUCCESS);
+        CV_OCL_CHECK(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, this));
     }
     if (asyncEvent)
-        clReleaseEvent(asyncEvent);
+        CV_OCL_DBG_CHECK(clReleaseEvent(asyncEvent));
     return retval == CL_SUCCESS;
 }
 
@@ -2420,19 +2491,20 @@ bool Kernel::runTask(bool sync, const Queue& q)
     cl_command_queue qq = getQueue(q);
     cl_event asyncEvent = 0;
     cl_int retval = clEnqueueTask(qq, p->handle, 0, 0, sync ? 0 : &asyncEvent);
-    if( sync || retval != CL_SUCCESS )
+    CV_OCL_DBG_CHECK_RESULT(retval, cv::format("clEnqueueTask('%s') sync=%s", p->name.c_str(), sync ? "true" : "false").c_str());
+    if (sync || retval != CL_SUCCESS)
     {
-        CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS);
+        CV_OCL_DBG_CHECK(clFinish(qq));
         p->cleanupUMats();
     }
     else
     {
         p->addref();
         p->isInProgress = true;
-        CV_OclDbgAssert(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS);
+        CV_OCL_CHECK(clSetEventCallback(asyncEvent, CL_COMPLETE, oclCleanupCallback, p));
     }
     if (asyncEvent)
-        clReleaseEvent(asyncEvent);
+        CV_OCL_DBG_CHECK(clReleaseEvent(asyncEvent));
     return retval == CL_SUCCESS;
 }
 
@@ -2454,8 +2526,9 @@ size_t Kernel::workGroupSize() const
         return 0;
     size_t val = 0, retsz = 0;
     cl_device_id dev = (cl_device_id)Device::getDefault().ptr();
-    return clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_WORK_GROUP_SIZE,
-                                    sizeof(val), &val, &retsz) == CL_SUCCESS ? val : 0;
+    cl_int status = clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_WORK_GROUP_SIZE, sizeof(val), &val, &retsz);
+    CV_OCL_CHECK_RESULT(status, "clGetKernelWorkGroupInfo(CL_KERNEL_WORK_GROUP_SIZE)");
+    return status == CL_SUCCESS ? val : 0;
 }
 
 size_t Kernel::preferedWorkGroupSizeMultiple() const
@@ -2464,8 +2537,9 @@ size_t Kernel::preferedWorkGroupSizeMultiple() const
         return 0;
     size_t val = 0, retsz = 0;
     cl_device_id dev = (cl_device_id)Device::getDefault().ptr();
-    return clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
-                                    sizeof(val), &val, &retsz) == CL_SUCCESS ? val : 0;
+    cl_int status = clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(val), &val, &retsz);
+    CV_OCL_CHECK_RESULT(status, "clGetKernelWorkGroupInfo(CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE)");
+    return status == CL_SUCCESS ? val : 0;
 }
 
 bool Kernel::compileWorkGroupSize(size_t wsz[]) const
@@ -2474,8 +2548,9 @@ bool Kernel::compileWorkGroupSize(size_t wsz[]) const
         return 0;
     size_t retsz = 0;
     cl_device_id dev = (cl_device_id)Device::getDefault().ptr();
-    return clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
-                                    sizeof(wsz[0])*3, wsz, &retsz) == CL_SUCCESS;
+    cl_int status = clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof(wsz[0])*3, wsz, &retsz);
+    CV_OCL_CHECK_RESULT(status, "clGetKernelWorkGroupInfo(CL_KERNEL_COMPILE_WORK_GROUP_SIZE)");
+    return status == CL_SUCCESS;
 }
 
 size_t Kernel::localMemSize() const
@@ -2485,8 +2560,9 @@ size_t Kernel::localMemSize() const
     size_t retsz = 0;
     cl_ulong val = 0;
     cl_device_id dev = (cl_device_id)Device::getDefault().ptr();
-    return clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_LOCAL_MEM_SIZE,
-                                    sizeof(val), &val, &retsz) == CL_SUCCESS ? (size_t)val : 0;
+    cl_int status = clGetKernelWorkGroupInfo(p->handle, dev, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(val), &val, &retsz);
+    CV_OCL_CHECK_RESULT(status, "clGetKernelWorkGroupInfo(CL_KERNEL_LOCAL_MEM_SIZE)");
+    return status == CL_SUCCESS ? (size_t)val : 0;
 }
 
 
@@ -2637,7 +2713,8 @@ struct Program::Impl
         cl_int retval = 0;
 
         handle = clCreateProgramWithSource((cl_context)ctx.ptr(), 1, &srcptr, &srclen, &retval);
-        CV_OclDbgAssert(handle && retval == CL_SUCCESS);
+        CV_OCL_DBG_CHECK_RESULT(retval, "clCreateProgramWithSource");
+        CV_Assert(handle || retval != CL_SUCCESS);
         if (handle && retval == CL_SUCCESS)
         {
             int i, n = (int)ctx.ndevices();
@@ -2693,7 +2770,7 @@ struct Program::Impl
                 // it would break CV_OPENCL_ALWAYS_SHOW_BUILD_LOG mode
                 if (retval != CL_SUCCESS && handle)
                 {
-                    clReleaseProgram(handle);
+                    CV_OCL_DBG_CHECK(clReleaseProgram(handle));
                     handle = NULL;
                 }
             }
@@ -2731,7 +2808,7 @@ struct Program::Impl
         cl_int binstatus = 0, retval = 0;
         handle = clCreateProgramWithBinary((cl_context)ctx.ptr(), 1, (cl_device_id*)&devid,
                                            &codelen, &bin, &binstatus, &retval);
-        CV_OclDbgAssert(retval == CL_SUCCESS);
+        CV_OCL_DBG_CHECK_RESULT(retval, "clCreateProgramWithBinary");
     }
 
     String store()
@@ -3081,8 +3158,7 @@ public:
         entry.capacity_ = alignSize(size, (int)_allocationGranularity(size));
         Context& ctx = Context::getDefault();
         cl_int retval = CL_SUCCESS;
-        entry.clBuffer_ = clCreateBuffer((cl_context)ctx.ptr(), CL_MEM_READ_WRITE|createFlags_, entry.capacity_, 0, &retval);
-        CV_Assert(retval == CL_SUCCESS);
+        CV_OCL_CHECK_(entry.clBuffer_ = clCreateBuffer((cl_context)ctx.ptr(), CL_MEM_READ_WRITE|createFlags_, entry.capacity_, 0, &retval), retval);
         CV_Assert(entry.clBuffer_ != NULL);
         if(retval == CL_SUCCESS)
         {
@@ -3099,7 +3175,7 @@ public:
         CV_Assert(entry.clBuffer_ != NULL);
         LOG_BUFFER_POOL("OpenCL release buffer: %p, %lld (0x%llx) bytes\n",
                 entry.clBuffer_, (long long)entry.capacity_, (long long)entry.capacity_);
-        clReleaseMemObject(entry.clBuffer_);
+        CV_OCL_DBG_CHECK(clReleaseMemObject(entry.clBuffer_));
     }
 };
 
@@ -3458,7 +3534,7 @@ public:
                         cl_int status = svmFns->fn_clEnqueueSVMMap(q, CL_TRUE, CL_MAP_WRITE,
                                 handle, u->size,
                                 0, NULL, NULL);
-                        CV_Assert(status == CL_SUCCESS);
+                        CV_OCL_CHECK_RESULT(status, "clEnqueueSVMMap()");
 
                     }
                     memcpy(handle, u->origdata, u->size);
@@ -3466,7 +3542,7 @@ public:
                     {
                         CV_OPENCL_SVM_TRACE_P("clEnqueueSVMUnmap: %p\n", handle);
                         cl_int status = svmFns->fn_clEnqueueSVMUnmap(q, handle, 0, NULL, NULL);
-                        CV_Assert(status == CL_SUCCESS);
+                        CV_OCL_CHECK_RESULT(status, "clEnqueueSVMUnmap()");
                     }
 
                     tempUMatFlags = UMatData::TEMP_UMAT | UMatData::TEMP_COPIED_UMAT;
@@ -3490,6 +3566,7 @@ public:
                     tempUMatFlags |= UMatData::TEMP_COPIED_UMAT;
                 }
             }
+            CV_OCL_DBG_CHECK_RESULT(retval, "clCreateBuffer()");
             if(!handle || retval != CL_SUCCESS)
                 return false;
             u->handle = handle;
@@ -3580,7 +3657,7 @@ public:
                             cl_int status = svmFns->fn_clEnqueueSVMMap(q, CL_FALSE, CL_MAP_READ,
                                     u->handle, u->size,
                                     0, NULL, NULL);
-                            CV_Assert(status == CL_SUCCESS);
+                            CV_OCL_CHECK_RESULT(status, "clEnqueueSVMMap()");
                         }
                         clFinish(q);
                         memcpy(u->origdata, u->handle, u->size);
@@ -3588,7 +3665,7 @@ public:
                         {
                             CV_OPENCL_SVM_TRACE_P("clEnqueueSVMUnmap: %p\n", u->handle);
                             cl_int status = svmFns->fn_clEnqueueSVMUnmap(q, u->handle, 0, NULL, NULL);
-                            CV_Assert(status == CL_SUCCESS);
+                            CV_OCL_CHECK_RESULT(status, "clEnqueueSVMUnmap()");
                         }
                     }
                     else
@@ -3604,8 +3681,8 @@ public:
                     if( u->tempCopiedUMat() )
                     {
                         AlignedDataPtr<false, true> alignedPtr(u->origdata, u->size, CV_OPENCL_DATA_PTR_ALIGNMENT);
-                        CV_OclDbgAssert(clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, 0,
-                                            u->size, alignedPtr.getAlignedPtr(), 0, 0, 0) == CL_SUCCESS);
+                        CV_OCL_CHECK(clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, 0,
+                                            u->size, alignedPtr.getAlignedPtr(), 0, 0, 0));
                     }
                     else
                     {
@@ -3617,14 +3694,14 @@ public:
                             void* data = clEnqueueMapBuffer(q, (cl_mem)u->handle, CL_TRUE,
                                 (CL_MAP_READ | CL_MAP_WRITE),
                                 0, u->size, 0, 0, 0, &retval);
-                            CV_Assert(retval == CL_SUCCESS);
+                            CV_OCL_CHECK_RESULT(retval, "clEnqueueMapBuffer()");
                             CV_Assert(u->origdata == data);
                             if (u->originalUMatData)
                             {
                                 CV_Assert(u->originalUMatData->data == data);
                             }
-                            CV_OclDbgAssert(clEnqueueUnmapMemObject(q, (cl_mem)u->handle, data, 0, 0, 0) == CL_SUCCESS);
-                            CV_OclDbgAssert(clFinish(q) == CL_SUCCESS);
+                            CV_OCL_CHECK(clEnqueueUnmapMemObject(q, (cl_mem)u->handle, data, 0, 0, 0));
+                            CV_OCL_DBG_CHECK(clFinish(q));
                         }
                     }
                 }
@@ -3650,7 +3727,7 @@ public:
             else
 #endif
             {
-                clReleaseMemObject((cl_mem)u->handle);
+                CV_OCL_DBG_CHECK(clReleaseMemObject((cl_mem)u->handle));
             }
             u->handle = 0;
             u->markDeviceCopyObsolete(true);
@@ -3698,7 +3775,7 @@ public:
                     {
                         CV_OPENCL_SVM_TRACE_P("clEnqueueSVMUnmap: %p\n", u->handle);
                         cl_int status = svmFns->fn_clEnqueueSVMUnmap(q, u->handle, 0, NULL, NULL);
-                        CV_Assert(status == CL_SUCCESS);
+                        CV_OCL_CHECK_RESULT(status, "clEnqueueSVMUnmap()");
                     }
                 }
                 bufferPoolSVM.release((void*)u->handle);
@@ -3706,7 +3783,7 @@ public:
 #endif
             else
             {
-                clReleaseMemObject((cl_mem)u->handle);
+                CV_OCL_DBG_CHECK(clReleaseMemObject((cl_mem)u->handle));
             }
             u->handle = 0;
             u->markDeviceCopyObsolete(true);
@@ -3747,7 +3824,7 @@ public:
                             cl_int status = svmFns->fn_clEnqueueSVMMap(q, CL_FALSE, CL_MAP_READ | CL_MAP_WRITE,
                                     u->handle, u->size,
                                     0, NULL, NULL);
-                            CV_Assert(status == CL_SUCCESS);
+                            CV_OCL_CHECK_RESULT(status, "clEnqueueSVMMap()");
                             u->allocatorFlags_ |= svm::OPENCL_SVM_BUFFER_MAP;
                         }
                     }
@@ -3767,6 +3844,7 @@ public:
                     u->data = (uchar*)clEnqueueMapBuffer(q, (cl_mem)u->handle, CL_TRUE,
                                                          (CL_MAP_READ | CL_MAP_WRITE),
                                                          0, u->size, 0, 0, 0, &retval);
+                    CV_OCL_DBG_CHECK_RESULT(retval, cv::format("clEnqueueMapBuffer(sz=%lld)", (int64)u->size).c_str());
                 }
                 if (u->data && retval == CL_SUCCESS)
                 {
@@ -3793,8 +3871,8 @@ public:
 #ifdef HAVE_OPENCL_SVM
             CV_DbgAssert((u->allocatorFlags_ & svm::OPENCL_SVM_BUFFER_MASK) == 0);
 #endif
-            CV_Assert( clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, 0,
-                                           u->size, alignedPtr.getAlignedPtr(), 0, 0, 0) == CL_SUCCESS );
+            CV_OCL_CHECK(clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE,
+                    0, u->size, alignedPtr.getAlignedPtr(), 0, 0, 0));
             u->markHostCopyObsolete(false);
         }
     }
@@ -3828,7 +3906,7 @@ public:
                         CV_OPENCL_SVM_TRACE_P("clEnqueueSVMUnmap: %p\n", u->handle);
                         cl_int status = svmFns->fn_clEnqueueSVMUnmap(q, u->handle,
                                 0, NULL, NULL);
-                        CV_Assert(status == CL_SUCCESS);
+                        CV_OCL_CHECK_RESULT(status, "clEnqueueSVMUnmap()");
                         clFinish(q);
                         u->allocatorFlags_ &= ~svm::OPENCL_SVM_BUFFER_MAP;
                     }
@@ -3843,12 +3921,11 @@ public:
             if (u->refcount == 0)
             {
                 CV_Assert(u->mapcount-- == 1);
-                CV_Assert((retval = clEnqueueUnmapMemObject(q,
-                          (cl_mem)u->handle, u->data, 0, 0, 0)) == CL_SUCCESS);
+                CV_OCL_CHECK(retval = clEnqueueUnmapMemObject(q, (cl_mem)u->handle, u->data, 0, 0, 0));
                 if (Device::getDefault().isAMD())
                 {
                     // required for multithreaded applications (see stitching test)
-                    CV_OclDbgAssert(clFinish(q) == CL_SUCCESS);
+                    CV_OCL_DBG_CHECK(clFinish(q));
                 }
                 u->markDeviceMemMapped(false);
                 u->data = 0;
@@ -3862,8 +3939,8 @@ public:
 #ifdef HAVE_OPENCL_SVM
             CV_DbgAssert((u->allocatorFlags_ & svm::OPENCL_SVM_BUFFER_MASK) == 0);
 #endif
-            CV_Assert( (retval = clEnqueueWriteBuffer(q, (cl_mem)u->handle, CL_TRUE, 0,
-                                u->size, alignedPtr.getAlignedPtr(), 0, 0, 0)) == CL_SUCCESS );
+            CV_OCL_CHECK(retval = clEnqueueWriteBuffer(q, (cl_mem)u->handle, CL_TRUE,
+                                0, u->size, alignedPtr.getAlignedPtr(), 0, 0, 0));
             u->markDeviceCopyObsolete(false);
             u->markHostCopyObsolete(true);
         }
@@ -3984,7 +4061,7 @@ public:
                 cl_int status = svmFns->fn_clEnqueueSVMMap(q, CL_FALSE, CL_MAP_READ,
                         u->handle, u->size,
                         0, NULL, NULL);
-                CV_Assert(status == CL_SUCCESS);
+                CV_OCL_CHECK_RESULT(status, "clEnqueueSVMMap()");
             }
             clFinish(q);
             if( iscontinuous )
@@ -4022,7 +4099,7 @@ public:
                 CV_OPENCL_SVM_TRACE_P("clEnqueueSVMUnmap: %p\n", u->handle);
                 cl_int status = svmFns->fn_clEnqueueSVMUnmap(q, u->handle,
                         0, NULL, NULL);
-                CV_Assert(status == CL_SUCCESS);
+                CV_OCL_CHECK_RESULT(status, "clEnqueueSVMUnmap()");
                 clFinish(q);
             }
         }
@@ -4032,19 +4109,19 @@ public:
             if( iscontinuous )
             {
                 AlignedDataPtr<false, true> alignedPtr((uchar*)dstptr, total, CV_OPENCL_DATA_PTR_ALIGNMENT);
-                CV_Assert(clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE,
-                    srcrawofs, total, alignedPtr.getAlignedPtr(), 0, 0, 0) >= 0 );
+                CV_OCL_CHECK(clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE,
+                    srcrawofs, total, alignedPtr.getAlignedPtr(), 0, 0, 0));
             }
             else
             {
                 AlignedDataPtr2D<false, true> alignedPtr((uchar*)dstptr, new_sz[1], new_sz[0], new_dststep[0], CV_OPENCL_DATA_PTR_ALIGNMENT);
                 uchar* ptr = alignedPtr.getAlignedPtr();
 
-                CV_Assert( clEnqueueReadBufferRect(q, (cl_mem)u->handle, CL_TRUE,
+                CV_OCL_CHECK(clEnqueueReadBufferRect(q, (cl_mem)u->handle, CL_TRUE,
                     new_srcofs, new_dstofs, new_sz,
                     new_srcstep[0], 0,
                     new_dststep[0], 0,
-                    ptr, 0, 0, 0) >= 0 );
+                    ptr, 0, 0, 0));
             }
         }
     }
@@ -4101,7 +4178,7 @@ public:
                 cl_int status = svmFns->fn_clEnqueueSVMMap(q, CL_FALSE, CL_MAP_WRITE,
                         u->handle, u->size,
                         0, NULL, NULL);
-                CV_Assert(status == CL_SUCCESS);
+                CV_OCL_CHECK_RESULT(status, "clEnqueueSVMMap()");
             }
             clFinish(q);
             if( iscontinuous )
@@ -4139,7 +4216,7 @@ public:
                 CV_OPENCL_SVM_TRACE_P("clEnqueueSVMUnmap: %p\n", u->handle);
                 cl_int status = svmFns->fn_clEnqueueSVMUnmap(q, u->handle,
                         0, NULL, NULL);
-                CV_Assert(status == CL_SUCCESS);
+                CV_OCL_CHECK_RESULT(status, "clEnqueueSVMUnmap()");
                 clFinish(q);
             }
         }
@@ -4149,19 +4226,19 @@ public:
             if( iscontinuous )
             {
                 AlignedDataPtr<true, false> alignedPtr((uchar*)srcptr, total, CV_OPENCL_DATA_PTR_ALIGNMENT);
-                CV_Assert(clEnqueueWriteBuffer(q, (cl_mem)u->handle, CL_TRUE,
-                    dstrawofs, total, alignedPtr.getAlignedPtr(), 0, 0, 0) >= 0);
+                CV_OCL_CHECK(clEnqueueWriteBuffer(q, (cl_mem)u->handle, CL_TRUE,
+                    dstrawofs, total, alignedPtr.getAlignedPtr(), 0, 0, 0));
             }
             else
             {
                 AlignedDataPtr2D<true, false> alignedPtr((uchar*)srcptr, new_sz[1], new_sz[0], new_srcstep[0], CV_OPENCL_DATA_PTR_ALIGNMENT);
                 uchar* ptr = alignedPtr.getAlignedPtr();
 
-                CV_Assert(clEnqueueWriteBufferRect(q, (cl_mem)u->handle, CL_TRUE,
+                CV_OCL_CHECK(clEnqueueWriteBufferRect(q, (cl_mem)u->handle, CL_TRUE,
                     new_dstofs, new_srcofs, new_sz,
                     new_dststep[0], 0,
                     new_srcstep[0], 0,
-                    ptr, 0, 0, 0) >= 0 );
+                    ptr, 0, 0, 0));
             }
         }
         u->markHostCopyObsolete(true);
@@ -4244,7 +4321,7 @@ public:
                     cl_int status = svmFns->fn_clEnqueueSVMMemcpy(q, CL_TRUE,
                             (uchar*)dst->handle + dstrawofs, (uchar*)src->handle + srcrawofs,
                             total, 0, NULL, NULL);
-                    CV_Assert(status == CL_SUCCESS);
+                    CV_OCL_CHECK_RESULT(status, "clEnqueueSVMMemcpy()");
                 }
                 else
                 {
@@ -4301,16 +4378,16 @@ public:
         {
             if( iscontinuous )
             {
-                CV_Assert( (retval = clEnqueueCopyBuffer(q, (cl_mem)src->handle, (cl_mem)dst->handle,
-                                               srcrawofs, dstrawofs, total, 0, 0, 0)) == CL_SUCCESS );
+                CV_OCL_CHECK(retval = clEnqueueCopyBuffer(q, (cl_mem)src->handle, (cl_mem)dst->handle,
+                                               srcrawofs, dstrawofs, total, 0, 0, 0));
             }
             else
             {
-                CV_Assert( (retval = clEnqueueCopyBufferRect(q, (cl_mem)src->handle, (cl_mem)dst->handle,
+                CV_OCL_CHECK(retval = clEnqueueCopyBufferRect(q, (cl_mem)src->handle, (cl_mem)dst->handle,
                                                    new_srcofs, new_dstofs, new_sz,
                                                    new_srcstep[0], 0,
                                                    new_dststep[0], 0,
-                                                   0, 0, 0)) == CL_SUCCESS );
+                                                   0, 0, 0));
             }
         }
         if (retval == CL_SUCCESS)
@@ -4333,7 +4410,7 @@ public:
 
         if( _sync )
         {
-            CV_OclDbgAssert(clFinish(q) == CL_SUCCESS);
+            CV_OCL_DBG_CHECK(clFinish(q));
         }
     }
 
@@ -4428,14 +4505,14 @@ void convertFromBuffer(void* cl_mem_buffer, size_t step, int rows, int cols, int
     cl_mem             memobj = (cl_mem)cl_mem_buffer;
     cl_mem_object_type mem_type = 0;
 
-    CV_Assert(clGetMemObjectInfo(memobj, CL_MEM_TYPE, sizeof(cl_mem_object_type), &mem_type, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetMemObjectInfo(memobj, CL_MEM_TYPE, sizeof(cl_mem_object_type), &mem_type, 0));
 
     CV_Assert(CL_MEM_OBJECT_BUFFER == mem_type);
 
     size_t total = 0;
-    CV_Assert(clGetMemObjectInfo(memobj, CL_MEM_SIZE, sizeof(size_t), &total, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetMemObjectInfo(memobj, CL_MEM_SIZE, sizeof(size_t), &total, 0));
 
-    CV_Assert(clRetainMemObject(memobj) == CL_SUCCESS);
+    CV_OCL_CHECK(clRetainMemObject(memobj));
 
     CV_Assert((int)step >= cols * CV_ELEM_SIZE(type));
     CV_Assert(total >= rows * step);
@@ -4465,12 +4542,12 @@ void convertFromImage(void* cl_mem_image, UMat& dst)
     cl_mem             clImage = (cl_mem)cl_mem_image;
     cl_mem_object_type mem_type = 0;
 
-    CV_Assert(clGetMemObjectInfo(clImage, CL_MEM_TYPE, sizeof(cl_mem_object_type), &mem_type, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetMemObjectInfo(clImage, CL_MEM_TYPE, sizeof(cl_mem_object_type), &mem_type, 0));
 
     CV_Assert(CL_MEM_OBJECT_IMAGE2D == mem_type);
 
     cl_image_format fmt = { 0, 0 };
-    CV_Assert(clGetImageInfo(clImage, CL_IMAGE_FORMAT, sizeof(cl_image_format), &fmt, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetImageInfo(clImage, CL_IMAGE_FORMAT, sizeof(cl_image_format), &fmt, 0));
 
     int depth = CV_8U;
     switch (fmt.image_channel_data_type)
@@ -4526,13 +4603,13 @@ void convertFromImage(void* cl_mem_image, UMat& dst)
     }
 
     size_t step = 0;
-    CV_Assert(clGetImageInfo(clImage, CL_IMAGE_ROW_PITCH, sizeof(size_t), &step, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetImageInfo(clImage, CL_IMAGE_ROW_PITCH, sizeof(size_t), &step, 0));
 
     size_t w = 0;
-    CV_Assert(clGetImageInfo(clImage, CL_IMAGE_WIDTH, sizeof(size_t), &w, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetImageInfo(clImage, CL_IMAGE_WIDTH, sizeof(size_t), &w, 0));
 
     size_t h = 0;
-    CV_Assert(clGetImageInfo(clImage, CL_IMAGE_HEIGHT, sizeof(size_t), &h, 0) == CL_SUCCESS);
+    CV_OCL_CHECK(clGetImageInfo(clImage, CL_IMAGE_HEIGHT, sizeof(size_t), &h, 0));
 
     dst.create((int)h, (int)w, type);
 
@@ -4543,9 +4620,9 @@ void convertFromImage(void* cl_mem_image, UMat& dst)
     size_t offset = 0;
     size_t src_origin[3] = { 0, 0, 0 };
     size_t region[3] = { w, h, 1 };
-    CV_Assert(clEnqueueCopyImageToBuffer(q, clImage, clBuffer, src_origin, region, offset, 0, NULL, NULL) == CL_SUCCESS);
+    CV_OCL_CHECK(clEnqueueCopyImageToBuffer(q, clImage, clBuffer, src_origin, region, offset, 0, NULL, NULL));
 
-    CV_Assert(clFinish(q) == CL_SUCCESS);
+    CV_OCL_CHECK(clFinish(q));
 
     return;
 } // convertFromImage()
@@ -4556,8 +4633,7 @@ void convertFromImage(void* cl_mem_image, UMat& dst)
 static void getDevices(std::vector<cl_device_id>& devices, cl_platform_id platform)
 {
     cl_uint numDevices = 0;
-    CV_OclDbgAssert(clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL,
-                                0, NULL, &numDevices) == CL_SUCCESS);
+    CV_OCL_DBG_CHECK(clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL, 0, NULL, &numDevices));
 
     if (numDevices == 0)
     {
@@ -4566,8 +4642,7 @@ static void getDevices(std::vector<cl_device_id>& devices, cl_platform_id platfo
     }
 
     devices.resize((size_t)numDevices);
-    CV_OclDbgAssert(clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL,
-                                numDevices, &devices[0], &numDevices) == CL_SUCCESS);
+    CV_OCL_DBG_CHECK(clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL, numDevices, &devices[0], &numDevices));
 }
 
 struct PlatformInfo::Impl
@@ -4658,7 +4733,7 @@ String PlatformInfo::version() const
 static void getPlatforms(std::vector<cl_platform_id>& platforms)
 {
     cl_uint numPlatforms = 0;
-    CV_OclDbgAssert(clGetPlatformIDs(0, NULL, &numPlatforms) == CL_SUCCESS);
+    CV_OCL_DBG_CHECK(clGetPlatformIDs(0, NULL, &numPlatforms));
 
     if (numPlatforms == 0)
     {
@@ -4667,7 +4742,7 @@ static void getPlatforms(std::vector<cl_platform_id>& platforms)
     }
 
     platforms.resize((size_t)numPlatforms);
-    CV_OclDbgAssert(clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms) == CL_SUCCESS);
+    CV_OCL_DBG_CHECK(clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms));
 }
 
 void getPlatfomsInfo(std::vector<PlatformInfo>& platformsInfo)
@@ -5048,11 +5123,12 @@ struct Image2D::Impl
         cl_int err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
                                                 CL_MEM_OBJECT_IMAGE2D, numFormats,
                                                 NULL, &numFormats);
+        CV_OCL_DBG_CHECK_RESULT(err, "clGetSupportedImageFormats(CL_MEM_OBJECT_IMAGE2D, NULL)");
         AutoBuffer<cl_image_format> formats(numFormats);
         err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
                                          CL_MEM_OBJECT_IMAGE2D, numFormats,
                                          formats, NULL);
-        CV_OclDbgAssert(err == CL_SUCCESS);
+        CV_OCL_DBG_CHECK_RESULT(err, "clGetSupportedImageFormats(CL_MEM_OBJECT_IMAGE2D, formats)");
         for (cl_uint i = 0; i < numFormats; ++i)
         {
             if (!memcmp(&formats[i], &format, sizeof(format)))
@@ -5113,7 +5189,7 @@ struct Image2D::Impl
             handle = clCreateImage2D(context, CL_MEM_READ_WRITE, &format, src.cols, src.rows, 0, NULL, &err);
             CV_SUPPRESS_DEPRECATED_END
         }
-        CV_OclDbgAssert(err == CL_SUCCESS);
+        CV_OCL_DBG_CHECK_RESULT(err, "clCreateImage()");
 
         size_t origin[] = { 0, 0, 0 };
         size_t region[] = { static_cast<size_t>(src.cols), static_cast<size_t>(src.rows), 1 };
@@ -5122,12 +5198,12 @@ struct Image2D::Impl
         if (!alias && !src.isContinuous())
         {
             devData = clCreateBuffer(context, CL_MEM_READ_ONLY, src.cols * src.rows * src.elemSize(), NULL, &err);
-            CV_OclDbgAssert(err == CL_SUCCESS);
+            CV_OCL_CHECK_RESULT(err, "clCreateBuffer()");
 
             const size_t roi[3] = {static_cast<size_t>(src.cols) * src.elemSize(), static_cast<size_t>(src.rows), 1};
-            CV_Assert(clEnqueueCopyBufferRect(queue, (cl_mem)src.handle(ACCESS_READ), devData, origin, origin,
-                roi, src.step, 0, src.cols * src.elemSize(), 0, 0, NULL, NULL) == CL_SUCCESS);
-            CV_OclDbgAssert(clFlush(queue) == CL_SUCCESS);
+            CV_OCL_CHECK(clEnqueueCopyBufferRect(queue, (cl_mem)src.handle(ACCESS_READ), devData, origin, origin,
+                roi, src.step, 0, src.cols * src.elemSize(), 0, 0, NULL, NULL));
+            CV_OCL_DBG_CHECK(clFlush(queue));
         }
         else
         {
@@ -5137,11 +5213,11 @@ struct Image2D::Impl
 
         if (!alias)
         {
-            CV_OclDbgAssert(clEnqueueCopyBufferToImage(queue, devData, handle, 0, origin, region, 0, NULL, 0) == CL_SUCCESS);
+            CV_OCL_CHECK(clEnqueueCopyBufferToImage(queue, devData, handle, 0, origin, region, 0, NULL, 0));
             if (!src.isContinuous())
             {
-                CV_OclDbgAssert(clFlush(queue) == CL_SUCCESS);
-                CV_OclDbgAssert(clReleaseMemObject(devData) == CL_SUCCESS);
+                CV_OCL_DBG_CHECK(clFlush(queue));
+                CV_OCL_DBG_CHECK(clReleaseMemObject(devData));
             }
         }
     }
@@ -5276,7 +5352,7 @@ struct Timer::Impl
     void start()
     {
 #ifdef HAVE_OPENCL
-        clFinish((cl_command_queue)queue.ptr());
+        CV_OCL_DBG_CHECK(clFinish((cl_command_queue)queue.ptr()));
         timer.start();
 #endif
     }
@@ -5284,7 +5360,7 @@ struct Timer::Impl
     void stop()
     {
 #ifdef HAVE_OPENCL
-        clFinish((cl_command_queue)queue.ptr());
+        CV_OCL_DBG_CHECK(clFinish((cl_command_queue)queue.ptr()));
         timer.stop();
 #endif
     }