diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index 8d9400224..3a28a3fdc 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -90,7 +90,8 @@ public: String vendor() const; String OpenCL_C_Version() const; String OpenCLVersion() const; - String deviceVersion() const; + int deviceVersionMajor() const; + int deviceVersionMinor() const; String driverVersion() const; void* ptr() const; @@ -224,16 +225,12 @@ public: static Context2& getDefault(bool initialize = true); void* ptr() const; - struct Impl; - inline struct Impl* _getImpl() const { return p; } + friend void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device); protected: + struct Impl; Impl* p; }; - -// TODO Move to internal header -void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device); - class CV_EXPORTS Platform { public: @@ -245,12 +242,14 @@ public: void* ptr() const; static Platform& getDefault(); - struct Impl; - inline struct Impl* _getImpl() const { return p; } + friend void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device); protected: + struct Impl; Impl* p; }; +// TODO Move to internal header +void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device); class CV_EXPORTS Queue { @@ -585,9 +584,12 @@ class CV_EXPORTS Image2D { public: Image2D(); - Image2D(const UMat &src); + explicit Image2D(const UMat &src); + Image2D(const Image2D & i); ~Image2D(); + Image2D & operator = (const Image2D & i); + void* ptr() const; protected: struct Impl; diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 37741c399..e64d09976 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -1505,6 +1505,9 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst) format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s%s", dcn, lcn, ocl::typeToStr(src.depth()), ocl::typeToStr(ddepth), doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut), ocl::KernelArg::WriteOnly(dst)); diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 510b17854..87e4fd57d 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -2915,6 +2915,9 @@ static bool ocl_transpose( InputArray _src, OutputArray _dst ) ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc, format("-D T=%s -D TILE_DIM=%d -D BLOCK_ROWS=%d", ocl::memopTypeToStr(type), TILE_DIM, BLOCK_ROWS)); + if (k.empty()) + return false; + if (inplace) k.args(ocl::KernelArg::ReadWriteNoSize(dst), dst.rows); else diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 8e6817b9f..d8254cbcb 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -1257,6 +1257,12 @@ OCL_FUNC(cl_int, clReleaseEvent, (cl_event event), (event)) #endif +#ifdef _DEBUG +#define CV_OclDbgAssert CV_DbgAssert +#else +#define CV_OclDbgAssert(expr) (void)(expr) +#endif + namespace cv { namespace ocl { struct UMat2D @@ -1539,6 +1545,8 @@ void finish2() void release() { if( CV_XADD(&refcount, -1) == 1 ) delete this; } \ int refcount +/////////////////////////////////////////// Platform ///////////////////////////////////////////// + struct Platform::Impl { Impl() @@ -1556,13 +1564,13 @@ struct Platform::Impl { //cl_uint num_entries cl_uint n = 0; - if( clGetPlatformIDs(1, &handle, &n) < 0 || n == 0 ) + if( clGetPlatformIDs(1, &handle, &n) != CL_SUCCESS || n == 0 ) handle = 0; if( handle != 0 ) { char buf[1000]; size_t len = 0; - clGetPlatformInfo(handle, CL_PLATFORM_VENDOR, sizeof(buf), buf, &len); + CV_OclDbgAssert(clGetPlatformInfo(handle, CL_PLATFORM_VENDOR, sizeof(buf), buf, &len) == CL_SUCCESS); buf[len] = '\0'; vendor = String(buf); } @@ -1623,7 +1631,29 @@ Platform& Platform::getDefault() return p; } -/////////////////////////////////////////////////////////////////////////////////// +/////////////////////////////////////// Device //////////////////////////////////////////// + +// deviceVersion has format +// OpenCL +// by specification +// http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html +// http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetDeviceInfo.html +static void parseDeviceVersion(const String &deviceVersion, int &major, int &minor) +{ + major = minor = 0; + if (10 >= deviceVersion.length()) + return; + const char *pstr = deviceVersion.c_str(); + if (0 != strncmp(pstr, "OpenCL ", 7)) + return; + size_t ppos = deviceVersion.find('.', 7); + if (String::npos == ppos) + return; + String temp = deviceVersion.substr(7, ppos - 7); + major = atoi(temp.c_str()); + temp = deviceVersion.substr(ppos + 1); + minor = atoi(temp.c_str()); +} struct Device::Impl { @@ -1639,8 +1669,10 @@ struct Device::Impl maxComputeUnits_ = getProp(CL_DEVICE_MAX_COMPUTE_UNITS); maxWorkGroupSize_ = getProp(CL_DEVICE_MAX_WORK_GROUP_SIZE); type_ = getProp(CL_DEVICE_TYPE); - deviceVersion_ = getStrProp(CL_DEVICE_VERSION); driverVersion_ = getStrProp(CL_DRIVER_VERSION); + + String deviceVersion_ = getStrProp(CL_DEVICE_VERSION); + parseDeviceVersion(deviceVersion_, deviceVersionMajor_, deviceVersionMinor_); } template @@ -1649,7 +1681,7 @@ struct Device::Impl _TpCL temp=_TpCL(); size_t sz = 0; - return clGetDeviceInfo(handle, prop, sizeof(temp), &temp, &sz) >= 0 && + return clGetDeviceInfo(handle, prop, sizeof(temp), &temp, &sz) == CL_SUCCESS && sz == sizeof(temp) ? _TpOut(temp) : _TpOut(); } @@ -1658,7 +1690,7 @@ struct Device::Impl cl_bool temp = CL_FALSE; size_t sz = 0; - return clGetDeviceInfo(handle, prop, sizeof(temp), &temp, &sz) >= 0 && + return clGetDeviceInfo(handle, prop, sizeof(temp), &temp, &sz) == CL_SUCCESS && sz == sizeof(temp) ? temp != 0 : false; } @@ -1666,7 +1698,7 @@ struct Device::Impl { char buf[1024]; size_t sz=0; - return clGetDeviceInfo(handle, prop, sizeof(buf)-16, buf, &sz) >= 0 && + return clGetDeviceInfo(handle, prop, sizeof(buf)-16, buf, &sz) == CL_SUCCESS && sz < sizeof(buf) ? String(buf) : String(); } @@ -1680,7 +1712,8 @@ struct Device::Impl int maxComputeUnits_; size_t maxWorkGroupSize_; int type_; - String deviceVersion_; + int deviceVersionMajor_; + int deviceVersionMinor_; String driverVersion_; }; @@ -1750,8 +1783,11 @@ String Device::OpenCL_C_Version() const String Device::OpenCLVersion() const { return p ? p->getStrProp(CL_DEVICE_EXTENSIONS) : String(); } -String Device::deviceVersion() const -{ return p ? p->deviceVersion_ : String(); } +int Device::deviceVersionMajor() const +{ return p ? p->deviceVersionMajor_ : 0; } + +int Device::deviceVersionMinor() const +{ return p ? p->deviceVersionMinor_ : 0; } String Device::driverVersion() const { return p ? p->driverVersion_ : String(); } @@ -1889,8 +1925,8 @@ void Device::maxWorkItemSizes(size_t* sizes) const { const int MAX_DIMS = 32; size_t retsz = 0; - clGetDeviceInfo(p->handle, CL_DEVICE_MAX_WORK_ITEM_SIZES, - MAX_DIMS*sizeof(sizes[0]), &sizes[0], &retsz); + CV_OclDbgAssert(clGetDeviceInfo(p->handle, CL_DEVICE_MAX_WORK_ITEM_SIZES, + MAX_DIMS*sizeof(sizes[0]), &sizes[0], &retsz) == CL_SUCCESS); } } @@ -1957,7 +1993,7 @@ const Device& Device::getDefault() return ctx.device(idx); } -///////////////////////////////////////////////////////////////////////////////////////// +////////////////////////////////////// Context /////////////////////////////////////////////////// template inline cl_int getStringInfo(Functor f, ObjectType obj, cl_uint name, std::string& param) @@ -1981,7 +2017,8 @@ inline cl_int getStringInfo(Functor f, ObjectType obj, cl_uint name, std::string return CL_SUCCESS; } -static void split(const std::string &s, char delim, std::vector &elems) { +static void split(const std::string &s, char delim, std::vector &elems) +{ elems.clear(); if (s.size() == 0) return; @@ -2023,15 +2060,12 @@ static bool parseOpenCLDeviceConfiguration(const std::string& configurationStr, static cl_device_id selectOpenCLDevice() { - std::string platform; + std::string platform, deviceName; std::vector deviceTypes; - std::string deviceName; + const char* configuration = getenv("OPENCV_OPENCL_DEVICE"); - if (configuration) - { - if (!parseOpenCLDeviceConfiguration(std::string(configuration), platform, deviceTypes, deviceName)) - return NULL; - } + if (configuration && !parseOpenCLDeviceConfiguration(std::string(configuration), platform, deviceTypes, deviceName)) + return NULL; bool isID = false; int deviceID = -1; @@ -2054,21 +2088,20 @@ static cl_device_id selectOpenCLDevice() if (isID) { deviceID = atoi(deviceName.c_str()); - CV_Assert(deviceID >= 0); + if (deviceID < 0) + return NULL; } } - cl_int status = CL_SUCCESS; std::vector platforms; { cl_uint numPlatforms = 0; - status = clGetPlatformIDs(0, NULL, &numPlatforms); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(clGetPlatformIDs(0, NULL, &numPlatforms) == CL_SUCCESS); + if (numPlatforms == 0) return NULL; platforms.resize((size_t)numPlatforms); - status = clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms) == CL_SUCCESS); platforms.resize(numPlatforms); } @@ -2078,8 +2111,7 @@ static cl_device_id selectOpenCLDevice() for (size_t i = 0; i < platforms.size(); i++) { std::string name; - status = getStringInfo(clGetPlatformInfo, platforms[i], CL_PLATFORM_NAME, name); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(getStringInfo(clGetPlatformInfo, platforms[i], CL_PLATFORM_NAME, name) == CL_SUCCESS); if (name.find(platform) != std::string::npos) { selectedPlatform = (int)i; @@ -2101,29 +2133,19 @@ static cl_device_id selectOpenCLDevice() deviceTypes.push_back("CPU"); } else - { deviceTypes.push_back("ALL"); - } } for (size_t t = 0; t < deviceTypes.size(); t++) { int deviceType = 0; if (deviceTypes[t] == "GPU") - { deviceType = Device::TYPE_GPU; - } else if (deviceTypes[t] == "CPU") - { deviceType = Device::TYPE_CPU; - } else if (deviceTypes[t] == "ACCELERATOR") - { deviceType = Device::TYPE_ACCELERATOR; - } else if (deviceTypes[t] == "ALL") - { deviceType = Device::TYPE_ALL; - } else { std::cerr << "ERROR: Unsupported device type for OpenCL device (GPU, CPU, ACCELERATOR): " << deviceTypes[t] << std::endl; @@ -2136,14 +2158,14 @@ static cl_device_id selectOpenCLDevice() i++) { cl_uint count = 0; - status = clGetDeviceIDs(platforms[i], deviceType, 0, NULL, &count); - CV_Assert(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND); + cl_int status = clGetDeviceIDs(platforms[i], deviceType, 0, NULL, &count); + CV_OclDbgAssert(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND); if (count == 0) continue; size_t base = devices.size(); devices.resize(base + count); status = clGetDeviceIDs(platforms[i], deviceType, count, &devices[base], &count); - CV_Assert(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND); + CV_OclDbgAssert(status == CL_SUCCESS || status == CL_DEVICE_NOT_FOUND); } for (size_t i = (isID ? deviceID : 0); @@ -2151,8 +2173,7 @@ static cl_device_id selectOpenCLDevice() i++) { std::string name; - status = getStringInfo(clGetDeviceInfo, devices[i], CL_DEVICE_NAME, name); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(getStringInfo(clGetDeviceInfo, devices[i], CL_DEVICE_NAME, name) == CL_SUCCESS); if (isID || name.find(deviceName) != std::string::npos) { // TODO check for OpenCL 1.1 @@ -2160,14 +2181,14 @@ static cl_device_id selectOpenCLDevice() } } } + not_found: std::cerr << "ERROR: Required OpenCL device not found, check configuration: " << (configuration == NULL ? "" : configuration) << std::endl << " Platform: " << (platform.length() == 0 ? "any" : platform) << std::endl << " Device types: "; for (size_t t = 0; t < deviceTypes.size(); t++) - { std::cerr << deviceTypes[t] << " "; - } + std::cerr << std::endl << " Device name: " << (deviceName.length() == 0 ? "any" : deviceName) << std::endl; return NULL; } @@ -2190,8 +2211,7 @@ struct Context2::Impl return; cl_platform_id pl = NULL; - cl_int status = clGetDeviceInfo(d, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &pl, NULL); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(clGetDeviceInfo(d, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &pl, NULL) == CL_SUCCESS); cl_context_properties prop[] = { @@ -2200,20 +2220,19 @@ struct Context2::Impl }; // !!! in the current implementation force the number of devices to 1 !!! - int nd = 1; + cl_uint nd = 1; + cl_int status; handle = clCreateContext(prop, nd, &d, 0, 0, &status); - CV_Assert(status == CL_SUCCESS); - bool ok = handle != 0 && status >= 0; + + bool ok = handle != 0 && status == CL_SUCCESS; if( ok ) { devices.resize(nd); devices[0].set(d); } else - { handle = NULL; - } } Impl(int dtype0) @@ -2231,13 +2250,12 @@ struct Context2::Impl cl_uint i, nd0 = 0, nd = 0; int dtype = dtype0 & 15; - clGetDeviceIDs( pl, dtype, 0, 0, &nd0 ); - if(retval < 0) - return; + CV_OclDbgAssert(clGetDeviceIDs( pl, dtype, 0, 0, &nd0 ) == CL_SUCCESS); + AutoBuffer dlistbuf(nd0*2+1); cl_device_id* dlist = (cl_device_id*)(void**)dlistbuf; cl_device_id* dlist_new = dlist + nd0; - clGetDeviceIDs( pl, dtype, nd0, dlist, &nd0 ); + CV_OclDbgAssert(clGetDeviceIDs( pl, dtype, nd0, dlist, &nd0 ) == CL_SUCCESS); String name0; for(i = 0; i < nd0; i++) @@ -2263,7 +2281,7 @@ struct Context2::Impl nd = 1; handle = clCreateContext(prop, nd, dlist_new, 0, 0, &retval); - bool ok = handle != 0 && retval >= 0; + bool ok = handle != 0 && retval == CL_SUCCESS; if( ok ) { devices.resize(nd); @@ -2275,7 +2293,10 @@ struct Context2::Impl ~Impl() { if(handle) + { clReleaseContext(handle); + handle = NULL; + } devices.clear(); } @@ -2431,11 +2452,10 @@ void initializeContextFromHandle(Context2& ctx, void* platform, void* _context, cl_device_id device = (cl_device_id)_device; // cleanup old context - Context2::Impl* impl = ctx._getImpl(); + Context2::Impl * impl = ctx.p; if (impl->handle) { - cl_int status = clReleaseContext(impl->handle); - (void)status; + CV_OclDbgAssert(clReleaseContext(impl->handle) == CL_SUCCESS); } impl->devices.clear(); @@ -2444,10 +2464,11 @@ void initializeContextFromHandle(Context2& ctx, void* platform, void* _context, impl->devices[0].set(device); Platform& p = Platform::getDefault(); - Platform::Impl* pImpl = p._getImpl(); + Platform::Impl* pImpl = p.p; pImpl->handle = (cl_platform_id)platform; } +/////////////////////////////////////////// Queue ///////////////////////////////////////////// struct Queue::Impl { @@ -2466,6 +2487,7 @@ struct Queue::Impl dh = (cl_device_id)pc->device(0).ptr(); cl_int retval = 0; handle = clCreateCommandQueue(ch, dh, 0, &retval); + CV_OclDbgAssert(retval == CL_SUCCESS); } ~Impl() @@ -2478,6 +2500,7 @@ struct Queue::Impl { clFinish(handle); clReleaseCommandQueue(handle); + handle = NULL; } } } @@ -2534,7 +2557,9 @@ bool Queue::create(const Context2& c, const Device& d) void Queue::finish() { if(p && p->handle) - clFinish(p->handle); + { + CV_OclDbgAssert(clFinish(p->handle) == CL_SUCCESS); + } } void* Queue::ptr() const @@ -2558,6 +2583,8 @@ static cl_command_queue getQueue(const Queue& q) return qq; } +/////////////////////////////////////////// KernelArg ///////////////////////////////////////////// + KernelArg::KernelArg() : flags(0), m(0), obj(0), sz(0), wscale(1) { @@ -2574,6 +2601,7 @@ KernelArg KernelArg::Constant(const Mat& m) return KernelArg(CONSTANT, 0, 1, m.data, m.total()*m.elemSize()); } +/////////////////////////////////////////// Kernel ///////////////////////////////////////////// struct Kernel::Impl { @@ -2584,6 +2612,7 @@ struct Kernel::Impl cl_int retval = 0; handle = ph != 0 ? clCreateKernel(ph, kname, &retval) : 0; + CV_OclDbgAssert(retval == CL_SUCCESS); for( int i = 0; i < MAX_ARRS; i++ ) u[i] = 0; haveTempDstUMats = false; @@ -2772,44 +2801,44 @@ int Kernel::set(int i, const KernelArg& arg) } if (ptronly) - clSetKernelArg(p->handle, (cl_uint)i++, sizeof(h), &h); + CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i++, sizeof(h), &h) == CL_SUCCESS); else if( arg.m->dims <= 2 ) { UMat2D u2d(*arg.m); - clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h); - clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u2d.step), &u2d.step); - clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u2d.offset), &u2d.offset); + CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h) == CL_SUCCESS); + 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); i += 3; if( !(arg.flags & KernelArg::NO_SIZE) ) { int cols = u2d.cols*arg.wscale; - clSetKernelArg(p->handle, (cl_uint)i, sizeof(u2d.rows), &u2d.rows); - clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(cols), &cols); + 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); i += 2; } } else { UMat3D u3d(*arg.m); - clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h); - clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u3d.slicestep), &u3d.slicestep); - clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u3d.step), &u3d.step); - clSetKernelArg(p->handle, (cl_uint)(i+3), sizeof(u3d.offset), &u3d.offset); + CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h) == CL_SUCCESS); + 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); i += 4; if( !(arg.flags & KernelArg::NO_SIZE) ) { int cols = u3d.cols*arg.wscale; - clSetKernelArg(p->handle, (cl_uint)i, sizeof(u3d.slices), &u3d.rows); - clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u3d.rows), &u3d.rows); - clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u3d.cols), &cols); + CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, sizeof(u3d.slices), &u3d.rows) == 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); i += 3; } } p->addUMat(*arg.m, (accessFlags & ACCESS_WRITE) != 0); return i; } - clSetKernelArg(p->handle, (cl_uint)i, arg.sz, arg.obj); + CV_OclDbgAssert(clSetKernelArg(p->handle, (cl_uint)i, arg.sz, arg.obj) == CL_SUCCESS); return i+1; } @@ -2839,17 +2868,17 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[], cl_int retval = clEnqueueNDRangeKernel(qq, p->handle, (cl_uint)dims, offset, globalsize, _localsize, 0, 0, sync ? 0 : &p->e); - if( sync || retval < 0 ) + if( sync || retval != CL_SUCCESS ) { - clFinish(qq); + CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS); p->cleanupUMats(); } else { p->addref(); - clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p); + CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS); } - return retval >= 0; + return retval == CL_SUCCESS; } bool Kernel::runTask(bool sync, const Queue& q) @@ -2859,62 +2888,62 @@ bool Kernel::runTask(bool sync, const Queue& q) cl_command_queue qq = getQueue(q); cl_int retval = clEnqueueTask(qq, p->handle, 0, 0, sync ? 0 : &p->e); - if( sync || retval < 0 ) + if( sync || retval != CL_SUCCESS ) { - clFinish(qq); + CV_OclDbgAssert(clFinish(qq) == CL_SUCCESS); p->cleanupUMats(); } else { p->addref(); - clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p); + CV_OclDbgAssert(clSetEventCallback(p->e, CL_COMPLETE, oclCleanupCallback, p) == CL_SUCCESS); } - return retval >= 0; + return retval == CL_SUCCESS; } size_t Kernel::workGroupSize() const { - if(!p) + if(!p || !p->handle) 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) >= 0 ? val : 0; + sizeof(val), &val, &retsz) == CL_SUCCESS ? val : 0; } size_t Kernel::preferedWorkGroupSizeMultiple() const { - if(!p) + if(!p || !p->handle) 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) >= 0 ? val : 0; + sizeof(val), &val, &retsz) == CL_SUCCESS ? val : 0; } bool Kernel::compileWorkGroupSize(size_t wsz[]) const { - if(!p || !wsz) + if(!p || !p->handle || !wsz) 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) >= 0; + sizeof(wsz[0]*3), wsz, &retsz) == CL_SUCCESS; } size_t Kernel::localMemSize() const { - if(!p) + if(!p || !p->handle) return 0; 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) >= 0 ? (size_t)val : 0; + sizeof(val), &val, &retsz) == CL_SUCCESS ? (size_t)val : 0; } -//////////////////////////////////////////////////////////////////////////////////////// +/////////////////////////////////////////// Program ///////////////////////////////////////////// struct Program::Impl { @@ -2931,7 +2960,7 @@ struct Program::Impl cl_int retval = 0; handle = clCreateProgramWithSource((cl_context)ctx.ptr(), 1, &srcptr, &srclen, &retval); - if( handle && retval >= 0 ) + if( handle && retval == CL_SUCCESS ) { int i, n = (int)ctx.ndevices(); AutoBuffer deviceListBuf(n+1); @@ -2942,21 +2971,22 @@ struct Program::Impl retval = clBuildProgram(handle, n, (const cl_device_id*)deviceList, buildflags.c_str(), 0, 0); - if( retval < 0 ) + if( retval != CL_SUCCESS ) { size_t retsz = 0; retval = clGetProgramBuildInfo(handle, (cl_device_id)deviceList[0], CL_PROGRAM_BUILD_LOG, 0, 0, &retsz); - if( retval >= 0 && retsz > 1 ) + if( retval == CL_SUCCESS && retsz > 1 ) { AutoBuffer bufbuf(retsz + 16); char* buf = bufbuf; retval = clGetProgramBuildInfo(handle, (cl_device_id)deviceList[0], CL_PROGRAM_BUILD_LOG, retsz+1, buf, &retsz); - if( retval >= 0 ) + if( retval == CL_SUCCESS ) { errmsg = String(buf); printf("OpenCL program can not be built: %s", errmsg.c_str()); + fflush(stdout); } } @@ -2999,6 +3029,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); } String store() @@ -3008,13 +3039,13 @@ struct Program::Impl size_t progsz = 0, retsz = 0; String prefix = Program::getPrefix(buildflags); size_t prefixlen = prefix.length(); - if(clGetProgramInfo(handle, CL_PROGRAM_BINARY_SIZES, sizeof(progsz), &progsz, &retsz) < 0) + if(clGetProgramInfo(handle, CL_PROGRAM_BINARY_SIZES, sizeof(progsz), &progsz, &retsz) != CL_SUCCESS) return String(); AutoBuffer bufbuf(prefixlen + progsz + 16); uchar* buf = bufbuf; memcpy(buf, prefix.c_str(), prefixlen); buf += prefixlen; - if(clGetProgramInfo(handle, CL_PROGRAM_BINARIES, sizeof(buf), &buf, &retsz) < 0) + if(clGetProgramInfo(handle, CL_PROGRAM_BINARIES, sizeof(buf), &buf, &retsz) != CL_SUCCESS) return String(); buf[progsz] = (uchar)'\0'; return String((const char*)(uchar*)bufbuf, prefixlen + progsz); @@ -3023,7 +3054,10 @@ struct Program::Impl ~Impl() { if( handle ) + { clReleaseProgram(handle); + handle = NULL; + } } IMPLEMENT_REFCOUNTABLE(); @@ -3123,7 +3157,7 @@ String Program::getPrefix(const String& buildflags) dev.name().c_str(), dev.driverVersion().c_str(), buildflags.c_str()); } -//////////////////////////////////////////////////////////////////////////////////////// +///////////////////////////////////////// ProgramSource2 /////////////////////////////////////////////// struct ProgramSource2::Impl { @@ -3198,7 +3232,7 @@ ProgramSource2::hash_t ProgramSource2::hash() const return p ? p->h : 0; } -////////////////////////////////////////////////////////////////////////////////////////////// +//////////////////////////////////////////// OpenCLAllocator ////////////////////////////////////////////////// class OpenCLAllocator : public MatAllocator { @@ -3243,7 +3277,7 @@ public: cl_int retval = 0; void* handle = clCreateBuffer((cl_context)ctx.ptr(), createFlags, total, 0, &retval); - if( !handle || retval < 0 ) + if( !handle || retval != CL_SUCCESS ) return defaultAllocate(dims, sizes, type, data, step, flags); UMatData* u = new UMatData(this); u->data = 0; @@ -3273,13 +3307,13 @@ public: int tempUMatFlags = UMatData::TEMP_UMAT; u->handle = clCreateBuffer(ctx_handle, CL_MEM_USE_HOST_PTR|createFlags, u->size, u->origdata, &retval); - if((!u->handle || retval < 0) && !(accessFlags & ACCESS_FAST)) + if((!u->handle || retval != CL_SUCCESS) && !(accessFlags & ACCESS_FAST)) { u->handle = clCreateBuffer(ctx_handle, CL_MEM_COPY_HOST_PTR|createFlags, u->size, u->origdata, &retval); tempUMatFlags = UMatData::TEMP_COPIED_UMAT; } - if(!u->handle || retval < 0) + if(!u->handle || retval != CL_SUCCESS) return false; u->prevAllocator = u->currAllocator; u->currAllocator = this; @@ -3339,8 +3373,8 @@ public: cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr(); if( u->tempCopiedUMat() ) { - clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, 0, - u->size, u->origdata, 0, 0, 0); + CV_OclDbgAssert(clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, 0, + u->size, u->origdata, 0, 0, 0) == CL_SUCCESS); } else { @@ -3348,8 +3382,9 @@ public: void* data = clEnqueueMapBuffer(q, (cl_mem)u->handle, CL_TRUE, (CL_MAP_READ | CL_MAP_WRITE), 0, u->size, 0, 0, 0, &retval); - clEnqueueUnmapMemObject(q, (cl_mem)u->handle, data, 0, 0, 0); - clFinish(q); + CV_OclDbgAssert(retval == CL_SUCCESS); + CV_OclDbgAssert(clEnqueueUnmapMemObject(q, (cl_mem)u->handle, data, 0, 0, 0) == CL_SUCCESS); + CV_OclDbgAssert(clFinish(q) == CL_SUCCESS); } } u->markHostCopyObsolete(false); @@ -3401,7 +3436,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); - if(u->data && retval >= 0) + if(u->data && retval == CL_SUCCESS) { u->markHostCopyObsolete(false); return; @@ -3421,7 +3456,7 @@ public: if( (accessFlags & ACCESS_READ) != 0 && u->hostCopyObsolete() ) { CV_Assert( clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, 0, - u->size, u->data, 0, 0, 0) >= 0 ); + u->size, u->data, 0, 0, 0) == CL_SUCCESS ); u->markHostCopyObsolete(false); } } @@ -3440,14 +3475,14 @@ public: if( !u->copyOnMap() && u->data ) { CV_Assert( (retval = clEnqueueUnmapMemObject(q, - (cl_mem)u->handle, u->data, 0, 0, 0)) >= 0 ); - clFinish(q); + (cl_mem)u->handle, u->data, 0, 0, 0)) == CL_SUCCESS ); + CV_OclDbgAssert(clFinish(q) == CL_SUCCESS); u->data = 0; } else if( u->copyOnMap() && u->deviceCopyObsolete() ) { CV_Assert( (retval = clEnqueueWriteBuffer(q, (cl_mem)u->handle, CL_TRUE, 0, - u->size, u->data, 0, 0, 0)) >= 0 ); + u->size, u->data, 0, 0, 0)) == CL_SUCCESS ); } u->markDeviceCopyObsolete(false); u->markHostCopyObsolete(false); @@ -3555,13 +3590,13 @@ public: if( iscontinuous ) { CV_Assert( clEnqueueReadBuffer(q, (cl_mem)u->handle, CL_TRUE, - srcrawofs, total, dstptr, 0, 0, 0) >= 0 ); + srcrawofs, total, dstptr, 0, 0, 0) == CL_SUCCESS ); } else { CV_Assert( clEnqueueReadBufferRect(q, (cl_mem)u->handle, CL_TRUE, new_srcofs, new_dstofs, new_sz, new_srcstep[0], new_srcstep[1], - new_dststep[0], new_dststep[1], dstptr, 0, 0, 0) >= 0 ); + new_dststep[0], new_dststep[1], dstptr, 0, 0, 0) == CL_SUCCESS ); } } @@ -3605,13 +3640,13 @@ public: if( iscontinuous ) { CV_Assert( clEnqueueWriteBuffer(q, (cl_mem)u->handle, - CL_TRUE, dstrawofs, total, srcptr, 0, 0, 0) >= 0 ); + CL_TRUE, dstrawofs, total, srcptr, 0, 0, 0) == CL_SUCCESS ); } else { CV_Assert( clEnqueueWriteBufferRect(q, (cl_mem)u->handle, CL_TRUE, new_dstofs, new_srcofs, new_sz, new_dststep[0], new_dststep[1], - new_srcstep[0], new_srcstep[1], srcptr, 0, 0, 0) >= 0 ); + new_srcstep[0], new_srcstep[1], srcptr, 0, 0, 0) == CL_SUCCESS ); } u->markHostCopyObsolete(true); @@ -3657,7 +3692,7 @@ public: if( iscontinuous ) { CV_Assert( clEnqueueCopyBuffer(q, (cl_mem)src->handle, (cl_mem)dst->handle, - srcrawofs, dstrawofs, total, 0, 0, 0) >= 0 ); + srcrawofs, dstrawofs, total, 0, 0, 0) == CL_SUCCESS ); } else { @@ -3666,14 +3701,16 @@ public: new_srcofs, new_dstofs, new_sz, new_srcstep[0], new_srcstep[1], new_dststep[0], new_dststep[1], - 0, 0, 0)) >= 0 ); + 0, 0, 0)) == CL_SUCCESS ); } dst->markHostCopyObsolete(true); dst->markDeviceCopyObsolete(false); if( _sync ) - clFinish(q); + { + CV_OclDbgAssert(clFinish(q) == CL_SUCCESS); + } } MatAllocator* matStdAllocator; @@ -3685,20 +3722,23 @@ MatAllocator* getOpenCLAllocator() return &allocator; } -/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// +///////////////////////////////////////////// Utility functions ///////////////////////////////////////////////// -static void getDevices(std::vector& devices,cl_platform_id& platform) +static void getDevices(std::vector& devices, cl_platform_id platform) { - cl_int status = CL_SUCCESS; cl_uint numDevices = 0; - status = clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL, 0, NULL, &numDevices); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL, + 0, NULL, &numDevices) == CL_SUCCESS); + if (numDevices == 0) + { + devices.clear(); return; + } + devices.resize((size_t)numDevices); - status = clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL, numDevices, &devices[0], &numDevices); - CV_Assert(status == CL_SUCCESS); - devices.resize(numDevices); + CV_OclDbgAssert(clGetDeviceIDs(platform, (cl_device_type)Device::TYPE_ALL, + numDevices, &devices[0], &numDevices) == CL_SUCCESS); } struct PlatformInfo2::Impl @@ -3714,7 +3754,7 @@ struct PlatformInfo2::Impl { char buf[1024]; size_t sz=0; - return clGetPlatformInfo(handle, prop, sizeof(buf)-16, buf, &sz) >= 0 && + return clGetPlatformInfo(handle, prop, sizeof(buf)-16, buf, &sz) == CL_SUCCESS && sz < sizeof(buf) ? String(buf) : String(); } @@ -3743,18 +3783,18 @@ PlatformInfo2::PlatformInfo2(const PlatformInfo2& i) { if (i.p) i.p->addref(); - this->p = i.p; + p = i.p; } PlatformInfo2& PlatformInfo2::operator =(const PlatformInfo2& i) { - if (i.p != this->p) + if (i.p != p) { if (i.p) i.p->addref(); - if (this->p) - this->p->release(); - this->p = i.p; + if (p) + p->release(); + p = i.p; } return *this; } @@ -3788,29 +3828,29 @@ String PlatformInfo2::version() const static void getPlatforms(std::vector& platforms) { - cl_int status = CL_SUCCESS; cl_uint numPlatforms = 0; - status = clGetPlatformIDs(0, NULL, &numPlatforms); - CV_Assert(status == CL_SUCCESS); + CV_OclDbgAssert(clGetPlatformIDs(0, NULL, &numPlatforms) == CL_SUCCESS); + if (numPlatforms == 0) + { + platforms.clear(); return; + } + platforms.resize((size_t)numPlatforms); - status = clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms); - CV_Assert(status == CL_SUCCESS); - platforms.resize(numPlatforms); + CV_OclDbgAssert(clGetPlatformIDs(numPlatforms, &platforms[0], &numPlatforms) == CL_SUCCESS); } void getPlatfomsInfo(std::vector& platformsInfo) { std::vector platforms; getPlatforms(platforms); + for (size_t i = 0; i < platforms.size(); i++) - { platformsInfo.push_back( PlatformInfo2((void*)&platforms[i]) ); - } } -const char* typeToStr(int t) +const char* typeToStr(int type) { static const char* tab[]= { @@ -3823,11 +3863,11 @@ const char* typeToStr(int t) "double", "double2", "double3", "double4", "?", "?", "?", "?" }; - int cn = CV_MAT_CN(t); - return cn > 4 ? "?" : tab[CV_MAT_DEPTH(t)*4 + cn-1]; + int cn = CV_MAT_CN(type), depth = CV_MAT_DEPTH(type); + return cn > 4 ? "?" : tab[depth*4 + cn-1]; } -const char* memopTypeToStr(int t) +const char* memopTypeToStr(int type) { static const char* tab[]= { @@ -3840,8 +3880,8 @@ const char* memopTypeToStr(int t) "int2", "int4", "?", "int8", "?", "?", "?", "?" }; - int cn = CV_MAT_CN(t); - return cn > 4 ? "?" : tab[CV_MAT_DEPTH(t)*4 + cn-1]; + int cn = CV_MAT_CN(type), depth = CV_MAT_DEPTH(type); + return cn > 4 ? "?" : tab[depth*4 + cn-1]; } const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf) @@ -3857,13 +3897,10 @@ const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf) sprintf(buf, "convert_%s", typestr); } else if( sdepth >= CV_32F ) - { sprintf(buf, "convert_%s%s_rte", typestr, (ddepth < CV_32S ? "_sat" : "")); - } else - { sprintf(buf, "convert_%s_sat", typestr); - } + return buf; } @@ -3919,28 +3956,7 @@ String kernelToStr(InputArray _kernel, int ddepth) return cv::format(" -D COEFF=%s", func(kernel).c_str()); } -/////////////////////////////////////////////////////////////////////////////////////////////// -// deviceVersion has format -// OpenCL -// by specification -// http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html -// http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetDeviceInfo.html -static void parseDeviceVersion(const String &deviceVersion, int &major, int &minor) -{ - major = minor = 0; - if (10 >= deviceVersion.length()) - return; - const char *pstr = deviceVersion.c_str(); - if (0 != strncmp(pstr, "OpenCL ", 7)) - return; - size_t ppos = deviceVersion.find('.', 7); - if (String::npos == ppos) - return; - String temp = deviceVersion.substr(7, ppos - 7); - major = atoi(temp.c_str()); - temp = deviceVersion.substr(ppos + 1); - minor = atoi(temp.c_str()); -} +/////////////////////////////////////////// Image2D //////////////////////////////////////////////////// struct Image2D::Impl { @@ -3950,54 +3966,41 @@ struct Image2D::Impl refcount = 1; init(src); } + ~Impl() { if (handle) clReleaseMemObject(handle); } + void init(const UMat &src) { - cl_image_format format; - int err; - int depth = src.depth(); - int channels = src.channels(); + CV_Assert(ocl::Device::getDefault().imageSupport()); + + cl_image_format format; + int err, depth = src.depth(), cn = src.channels(); + CV_Assert(cn <= 4); + + static const int channelTypes[] = { CL_UNSIGNED_INT8, CL_SIGNED_INT8, CL_UNSIGNED_INT16, + CL_SIGNED_INT16, CL_SIGNED_INT32, CL_FLOAT, -1, -1 }; + static const int channelOrders[] = { -1, CL_R, CL_RG, -1, CL_RGBA }; + + int channelType = channelTypes[depth], channelOrder = channelOrders[cn]; + if (channelType < 0 || channelOrder < 0) + CV_Error(Error::OpenCLApiCallError, "Image format is not supported"); + + format.image_channel_data_type = (cl_channel_type)channelType; + format.image_channel_order = (cl_channel_order)channelOrder; + + cl_context context = (cl_context)Context2::getDefault().ptr(); + cl_command_queue queue = (cl_command_queue)Queue::getDefault().ptr(); - switch(depth) - { - case CV_8U: - format.image_channel_data_type = CL_UNSIGNED_INT8; - break; - case CV_32S: - format.image_channel_data_type = CL_UNSIGNED_INT32; - break; - case CV_32F: - format.image_channel_data_type = CL_FLOAT; - break; - default: - CV_Error(-1, "Image forma is not supported"); - break; - } - switch(channels) - { - case 1: - format.image_channel_order = CL_R; - break; - case 3: - format.image_channel_order = CL_RGB; - break; - case 4: - format.image_channel_order = CL_RGBA; - break; - default: - CV_Error(-1, "Image format is not supported"); - break; - } #ifdef CL_VERSION_1_2 - //this enables backwards portability to - //run on OpenCL 1.1 platform if library binaries are compiled with OpenCL 1.2 support - int minor, major; - parseDeviceVersion(Device::getDefault().deviceVersion(), major, minor); - if ((1 < major) || ((1 == major) && (2 <= minor))) + // this enables backwards portability to + // run on OpenCL 1.1 platform if library binaries are compiled with OpenCL 1.2 support + const Device & d = ocl::Device::getDefault(); + int minor = d.deviceVersionMinor(), major = d.deviceVersionMajor(); + if (1 < major || (1 == major && 2 <= minor)) { cl_image_desc desc; desc.image_type = CL_MEM_OBJECT_IMAGE2D; @@ -4010,35 +4013,38 @@ struct Image2D::Impl desc.buffer = NULL; desc.num_mip_levels = 0; desc.num_samples = 0; - handle = clCreateImage((cl_context)Context2::getDefault().ptr(), CL_MEM_READ_WRITE, &format, &desc, NULL, &err); + handle = clCreateImage(context, CL_MEM_READ_WRITE, &format, &desc, NULL, &err); } else #endif { - handle = clCreateImage2D((cl_context)Context2::getDefault().ptr(), CL_MEM_READ_WRITE, &format, src.cols, src.rows, 0, NULL, &err); + handle = clCreateImage2D(context, CL_MEM_READ_WRITE, &format, src.cols, src.rows, 0, NULL, &err); } + CV_OclDbgAssert(err == CL_SUCCESS); + size_t origin[] = { 0, 0, 0 }; size_t region[] = { src.cols, src.rows, 1 }; cl_mem devData; if (!src.isContinuous()) { - devData = clCreateBuffer((cl_context)Context2::getDefault().ptr(), CL_MEM_READ_ONLY, src.cols * src.rows * src.elemSize(), NULL, NULL); + devData = clCreateBuffer(context, CL_MEM_READ_ONLY, src.cols * src.rows * src.elemSize(), NULL, &err); + CV_OclDbgAssert(err == CL_SUCCESS); + const size_t roi[3] = {src.cols * src.elemSize(), src.rows, 1}; - clEnqueueCopyBufferRect((cl_command_queue)Queue::getDefault().ptr(), (cl_mem)src.handle(ACCESS_READ), devData, origin, origin, - roi, src.step, 0, src.cols * src.elemSize(), 0, 0, NULL, NULL); - clFlush((cl_command_queue)Queue::getDefault().ptr()); + 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); } else - { devData = (cl_mem)src.handle(ACCESS_READ); - } + CV_Assert(devData != NULL); - clEnqueueCopyBufferToImage((cl_command_queue)Queue::getDefault().ptr(), devData, handle, 0, origin, region, 0, NULL, 0); + CV_OclDbgAssert(clEnqueueCopyBufferToImage(queue, devData, handle, 0, origin, region, 0, NULL, 0) == CL_SUCCESS); if (!src.isContinuous()) { - clFlush((cl_command_queue)Queue::getDefault().ptr()); - clReleaseMemObject(devData); + CV_OclDbgAssert(clFlush(queue) == CL_SUCCESS); + CV_OclDbgAssert(clReleaseMemObject(devData) == CL_SUCCESS); } } @@ -4051,10 +4057,32 @@ Image2D::Image2D() { p = NULL; } + Image2D::Image2D(const UMat &src) { p = new Impl(src); } + +Image2D::Image2D(const Image2D & i) +{ + p = i.p; + if (p) + p->addref(); +} + +Image2D & Image2D::operator = (const Image2D & i) +{ + if (i.p != p) + { + if (i.p) + i.p->addref(); + if (p) + p->release(); + p = i.p; + } + return *this; +} + Image2D::~Image2D() { if (p) diff --git a/modules/ts/src/ocl_test.cpp b/modules/ts/src/ocl_test.cpp index 0ad3df693..389b2aefc 100644 --- a/modules/ts/src/ocl_test.cpp +++ b/modules/ts/src/ocl_test.cpp @@ -160,17 +160,10 @@ void dumpOpenCLDevice() DUMP_MESSAGE_STDOUT(" Max memory allocation size = "<< maxMemAllocSizeStr); DUMP_PROPERTY_XML("cv_ocl_current_maxMemAllocSize", device.maxMemAllocSize()); -#if 0 - const char* doubleSupportStr = device.haveDoubleSupport() ? "Yes" : "No"; - DUMP_MESSAGE_STDOUT(" Double support = "<< doubleSupportStr); - DUMP_PROPERTY_XML("cv_ocl_current_haveDoubleSupport", device.haveDoubleSupport()); -#else const char* doubleSupportStr = device.doubleFPConfig() > 0 ? "Yes" : "No"; DUMP_MESSAGE_STDOUT(" Double support = "<< doubleSupportStr); DUMP_PROPERTY_XML("cv_ocl_current_haveDoubleSupport", device.doubleFPConfig() > 0); -#endif - const char* isUnifiedMemoryStr = device.hostUnifiedMemory() ? "Yes" : "No"; DUMP_MESSAGE_STDOUT(" Host unified memory = "<< isUnifiedMemoryStr); DUMP_PROPERTY_XML("cv_ocl_current_hostUnifiedMemory", device.hostUnifiedMemory()); diff --git a/modules/video/src/opencl/optical_flow_farneback.cl b/modules/video/src/opencl/optical_flow_farneback.cl index 0ef48d2c0..778583943 100644 --- a/modules/video/src/opencl/optical_flow_farneback.cl +++ b/modules/video/src/opencl/optical_flow_farneback.cl @@ -142,11 +142,6 @@ inline int idx_row_high(const int y, const int last_row) return abs(last_row - abs(last_row - y)) % (last_row + 1); } -inline int idx_row(const int y, const int last_row) -{ - return idx_row_low(idx_row_high(y, last_row), last_row); -} - inline int idx_col_low(const int x, const int last_col) { return abs(x) % (last_col + 1); @@ -431,4 +426,4 @@ __kernel void updateFlow(__global const float * M, int mStep, flowx[mad24(y, xStep, x)] = (g11*h2 - g12*h1) * detInv; flowy[mad24(y, yStep, x)] = (g22*h1 - g12*h2) * detInv; } -} \ No newline at end of file +}