generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Rewrite queryDeviceInfo interface.

Browse Source 
Previously the function may cause some unsafe issue. It is fixed now by introducing a template parameter.
This commit is contained in:
peng xiao 2013-05-09 17:15:26 +08:00
parent 416fb50594
commit bfa0f02303
2 changed files with 45 additions and 61 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
modules/ocl/include/opencv2/ocl/private/util.hpp
View File

@ -125,16 +125,21 @@ namespace cv
bool CV_EXPORTS support_image2d(Context *clCxt = Context::getContext());
// the enums are used to query device information
// currently only support wavefront size queries
enum DEVICE_INFO
{
WAVEFRONT_SIZE, //in AMD speak
WARP_SIZE = WAVEFRONT_SIZE, //in nvidia speak
IS_CPU_DEVICE //check if the device is CPU
WAVEFRONT_SIZE,
IS_CPU_DEVICE //supports
};
//info should have been pre-allocated
void CV_EXPORTS queryDeviceInfo(DEVICE_INFO info_type, void* info);
template<DEVICE_INFO _it, typename _ty>
_ty queryDeviceInfo(cl_kernel kernel = NULL);
//only these three specializations are implemented at the moment
template<>
int CV_EXPORTS queryDeviceInfo<WAVEFRONT_SIZE, int>(cl_kernel kernel);
template<>
size_t CV_EXPORTS queryDeviceInfo<WAVEFRONT_SIZE, size_t>(cl_kernel kernel);
template<>
bool CV_EXPORTS queryDeviceInfo<IS_CPU_DEVICE, bool>(cl_kernel kernel);
}//namespace ocl
}//namespace cv

89
modules/ocl/src/initialization.cpp
View File

@ -355,64 +355,43 @@ namespace cv
clFinish(Context::getContext()->impl->clCmdQueue);
}
void queryDeviceInfo(DEVICE_INFO info_type, void* info)
//template specializations of queryDeviceInfo
template<>
bool queryDeviceInfo<IS_CPU_DEVICE, bool>(cl_kernel)
{
static Info::Impl* impl = Context::getContext()->impl;
switch(info_type)
{
case WAVEFRONT_SIZE:
{
bool is_cpu = false;
queryDeviceInfo(IS_CPU_DEVICE, &is_cpu);
if(is_cpu)
{
*(int*)info = 1;
return;
}
#ifdef CL_DEVICE_WAVEFRONT_WIDTH_AMD
try
{
openCLSafeCall(clGetDeviceInfo(Context::getContext()->impl->devices[0],
CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof(size_t), info, 0));
}
catch(const cv::Exception&)
#elif defined (CL_DEVICE_WARP_SIZE_NV)
const int EXT_LEN = 4096 + 1 ;
char extends_set[EXT_LEN];
size_t extends_size;
openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum], CL_DEVICE_EXTENSIONS, EXT_LEN, (void *)extends_set, &extends_size));
extends_set[EXT_LEN - 1] = 0;
if(std::string(extends_set).find("cl_nv_device_attribute_query") != std::string::npos)
{
openCLSafeCall(clGetDeviceInfo(Context::getContext()->impl->devices[0],
CL_DEVICE_WARP_SIZE_NV, sizeof(size_t), info, 0));
}
else
#endif
{
// if no way left for us to query the warp size, we can get it from kernel group info
static const char * _kernel_string = "__kernel void test_func() {}";
cl_kernel kernel;
kernel = openCLGetKernelFromSource(Context::getContext(), &_kernel_string, "test_func");
openCLSafeCall(clGetKernelWorkGroupInfo(kernel, impl->devices[impl->devnum],
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(size_t), info, NULL));
}
Info::Impl* impl = Context::getContext()->impl;
cl_device_type devicetype;
openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum],
CL_DEVICE_TYPE, sizeof(cl_device_type),
&devicetype, NULL));
return (devicetype == CVCL_DEVICE_TYPE_CPU);
}
}
break;
case IS_CPU_DEVICE:
{
cl_device_type devicetype;
openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum],
CL_DEVICE_TYPE, sizeof(cl_device_type),
&devicetype, NULL));
*(bool*)info = (devicetype == CVCL_DEVICE_TYPE_CPU);
}
break;
default:
CV_Error(-1, "Invalid device info type");
break;
template<typename _ty>
static _ty queryWavesize(cl_kernel kernel)
{
size_t info = 0;
Info::Impl* impl = Context::getContext()->impl;
bool is_cpu = queryDeviceInfo<IS_CPU_DEVICE, bool>();
if(is_cpu)
{
return 1;
}
CV_Assert(kernel != NULL);
openCLSafeCall(clGetKernelWorkGroupInfo(kernel, impl->devices[impl->devnum],
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(size_t), &info, NULL));
return static_cast<_ty>(info);
}
template<>
size_t queryDeviceInfo<WAVEFRONT_SIZE, size_t>(cl_kernel kernel)
{
return queryWavesize<size_t>(kernel);
}
template<>
int queryDeviceInfo<WAVEFRONT_SIZE, int>(cl_kernel kernel)
{
return queryWavesize<int>(kernel);
}
void openCLReadBuffer(Context *clCxt, cl_mem dst_buffer, void *host_buffer, size_t size)