#endif
+#ifdef _DEBUG
+#define CV_OclDbgAssert CV_DbgAssert
+#else
+#define CV_OclDbgAssert(expr) (void)(expr)
+#endif
+
namespace cv { namespace ocl {
struct UMat2D
void release() { if( CV_XADD(&refcount, -1) == 1 ) delete this; } \
int refcount
+/////////////////////////////////////////// Platform /////////////////////////////////////////////
+
struct Platform::Impl
{
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);
}
return p;
}
-///////////////////////////////////////////////////////////////////////////////////
+/////////////////////////////////////// Device ////////////////////////////////////////////
+
+// deviceVersion has format
+// OpenCL<space><major_version.minor_version><space><vendor-specific information>
+// 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
{
maxComputeUnits_ = getProp<cl_uint, int>(CL_DEVICE_MAX_COMPUTE_UNITS);
maxWorkGroupSize_ = getProp<size_t, size_t>(CL_DEVICE_MAX_WORK_GROUP_SIZE);
type_ = getProp<cl_device_type, int>(CL_DEVICE_TYPE);
- deviceVersion_ = getStrProp(CL_DEVICE_VERSION);
driverVersion_ = getStrProp(CL_DRIVER_VERSION);
+
+ String deviceVersion_ = getStrProp(CL_DEVICE_VERSION);
+ parseDeviceVersion(deviceVersion_, deviceVersionMajor_, deviceVersionMinor_);
}
template<typename _TpCL, typename _TpOut>
_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();
}
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;
}
{
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();
}
int maxComputeUnits_;
size_t maxWorkGroupSize_;
int type_;
- String deviceVersion_;
+ int deviceVersionMajor_;
+ int deviceVersionMinor_;
String driverVersion_;
};
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(); }
{
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);
}
}
return ctx.device(idx);
}
-/////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////// Context ///////////////////////////////////////////////////
template <typename Functor, typename ObjectType>
inline cl_int getStringInfo(Functor f, ObjectType obj, cl_uint name, std::string& param)
return CL_SUCCESS;
}
-static void split(const std::string &s, char delim, std::vector<std::string> &elems) {
+static void split(const std::string &s, char delim, std::vector<std::string> &elems)
+{
elems.clear();
if (s.size() == 0)
return;
static cl_device_id selectOpenCLDevice()
{
- std::string platform;
+ std::string platform, deviceName;
std::vector<std::string> 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;
if (isID)
{
deviceID = atoi(deviceName.c_str());
- CV_Assert(deviceID >= 0);
+ if (deviceID < 0)
+ return NULL;
}
}
- cl_int status = CL_SUCCESS;
std::vector<cl_platform_id> 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);
}
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;
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;
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);
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
}
}
}
+
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;
}
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[] =
{
};
// !!! 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)
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<void*> 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++)
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);
~Impl()
{
if(handle)
+ {
clReleaseContext(handle);
+ handle = NULL;
+ }
devices.clear();
}
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();
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
{
dh = (cl_device_id)pc->device(0).ptr();
cl_int retval = 0;
handle = clCreateCommandQueue(ch, dh, 0, &retval);
+ CV_OclDbgAssert(retval == CL_SUCCESS);
}
~Impl()
{
clFinish(handle);
clReleaseCommandQueue(handle);
+ handle = NULL;
}
}
}
void Queue::finish()
{
if(p && p->handle)
- clFinish(p->handle);
+ {
+ CV_OclDbgAssert(clFinish(p->handle) == CL_SUCCESS);
+ }
}
void* Queue::ptr() const
return qq;
}
+/////////////////////////////////////////// KernelArg /////////////////////////////////////////////
+
KernelArg::KernelArg()
: flags(0), m(0), obj(0), sz(0), wscale(1)
{
return KernelArg(CONSTANT, 0, 1, m.data, m.total()*m.elemSize());
}
+/////////////////////////////////////////// Kernel /////////////////////////////////////////////
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;
}
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;
}
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)
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
{
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<void*> deviceListBuf(n+1);
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<char> 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);
}
}
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()
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<uchar> 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);
~Impl()
{
if( handle )
+ {
clReleaseProgram(handle);
+ handle = NULL;
+ }
}
IMPLEMENT_REFCOUNTABLE();
dev.name().c_str(), dev.driverVersion().c_str(), buildflags.c_str());
}
-////////////////////////////////////////////////////////////////////////////////////////
+///////////////////////////////////////// ProgramSource2 ///////////////////////////////////////////////
struct ProgramSource2::Impl
{
return p ? p->h : 0;
}
-//////////////////////////////////////////////////////////////////////////////////////////////
+//////////////////////////////////////////// OpenCLAllocator //////////////////////////////////////////////////
class OpenCLAllocator : public MatAllocator
{
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;
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;
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
{
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);
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;
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);
}
}
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);
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 );
}
}
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);
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
{
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;
return &allocator;
}
-///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+///////////////////////////////////////////// Utility functions /////////////////////////////////////////////////
-static void getDevices(std::vector<cl_device_id>& devices,cl_platform_id& platform)
+static void getDevices(std::vector<cl_device_id>& 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
{
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();
}
{
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;
}
static void getPlatforms(std::vector<cl_platform_id>& 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<PlatformInfo2>& platformsInfo)
{
std::vector<cl_platform_id> 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[]=
{
"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[]=
{
"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)
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;
}
return cv::format(" -D COEFF=%s", func(kernel).c_str());
}
-///////////////////////////////////////////////////////////////////////////////////////////////
-// deviceVersion has format
-// OpenCL<space><major_version.minor_version><space><vendor-specific information>
-// 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
{
refcount = 1;
init(src);
}
+
~Impl()
{
if (handle)
clReleaseMemObject(handle);
}
+
void init(const UMat &src)
{
+ CV_Assert(ocl::Device::getDefault().imageSupport());
+
cl_image_format format;
- int err;
- int depth = src.depth();
- int channels = src.channels();
+ 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;
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);
}
}
{
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)
projects / profile / ivi / opencv.git / commitdiff