diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index c50af047ad..5d7981aeb4 100644 --- a/modules/core/src/ocl.cpp +++ b/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 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 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 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 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); + 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 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 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 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 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 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) @@ -4517,7 +4594,7 @@ void convertFromImage(void* cl_mem_image, UMat& dst) case CL_RGBA: case CL_BGRA: case CL_ARGB: - type = CV_MAKE_TYPE(depth, 4); + type = CV_MAKE_TYPE(depth, 4); break; default: @@ -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& 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& 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& 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& 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& 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 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(src.cols), static_cast(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(src.cols) * src.elemSize(), static_cast(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 }