Skip to content

O
opencv
Commit 8fb48c09 authored by Alexander Alekhin's avatar Alexander Alekhin
Browse files
Options

ocl: improve debug information

parent 6e4f9433
Hide whitespace changes
Inline Side-by-side
Showing with 244 additions and 165 deletions
modules/core/src/ocl.cpp
View file @ 8fb48c09
......@@ -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,6 +97,14 @@
#include "ocl_deprecated.hpp"
#endif // HAVE_OPENCL
#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 {
#ifdef _DEBUG
#define CV_OclDbgAssert CV_DbgAssert
#else
......@@ -111,13 +122,53 @@ static bool isRaiseError()
#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 +479,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 +907,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 +1093,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 +1108,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 +1159,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 +1179,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 +1314,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 +1327,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 +1353,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 +1384,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 +1398,7 @@ struct Context::Impl
{
if(handle)
{
clReleaseContext(handle);
CV_OCL_DBG_CHECK(clReleaseContext(handle));
handle = NULL;
}
devices.clear();
......@@ -1527,8 +1586,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 +1806,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 +1825,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 +1864,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 +1884,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 +1914,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 +1941,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 +1953,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 +1969,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 +2040,7 @@ void Queue::finish()
{
if(p && p->handle)
{
CV_OclDbgAssert(clFinish(p->handle) == CL_SUCCESS);
CV_OCL_DBG_CHECK(clFinish(p->handle));
}
}
......@@ -2044,16 +2095,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 +2144,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 +2156,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 +2278,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 +2304,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 +2327,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 +2347,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 +2421,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 +2435,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 +2479,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 +2494,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 +2529,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 +2540,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 +2551,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 +2563,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,6 +2716,7 @@ struct Program::Impl
cl_int retval = 0;
handle = clCreateProgramWithSource((cl_context)ctx.ptr(), 1, &srcptr, &srclen, &retval);
CV_OCL_DBG_CHECK_RESULT(retval, "clCreateProgramWithSource");
CV_OclDbgAssert(handle && retval == CL_SUCCESS);
if (handle && retval == CL_SUCCESS)
{
......@@ -2693,7 +2773,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 +2811,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 +3161,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 +3178,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 +3537,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 +3545,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 +3569,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 +3660,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 +3668,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 +3684,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 +3697,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 +3730,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 +3778,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 +3786,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 +3827,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 +3847,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 +3874,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 +3909,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 +3924,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 +3942,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 +4064,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 +4102,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 +4112,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 +4181,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 +4219,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 +4229,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 +4324,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 +4381,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 +4413,7 @@ public:
if( _sync )
{
CV_OclDbgAssert(clFinish(q) == CL_SUCCESS);
CV_OCL_DBG_CHECK(clFinish(q));
}
}
......@@ -4428,14 +4508,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 +4545,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 +4597,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 +4606,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 +4623,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 +4636,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 +4645,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 +4736,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 +4745,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 +5126,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 +5192,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 +5201,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 +5216,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 +5355,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 +5363,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
}
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment