Copy ocl::queryDeviceInfo interface from master to 2.4.

Affected functions surf.ocl, pyrlk.ocl and hog.ocl are updated with the change.
This commit is contained in:
peng xiao 2013-05-14 17:50:38 +08:00
parent 2aa5f1bfd1
commit ac21cabda2
5 changed files with 62 additions and 92 deletions

View File

@ -60,27 +60,24 @@ namespace cv
const char noImage2dOption [] = "-D DISABLE_IMAGE2D"; const char noImage2dOption [] = "-D DISABLE_IMAGE2D";
static char SURF_OPTIONS [1024] = ""; static bool use_image2d = false;
static bool USE_IMAGE2d = false;
static void openCLExecuteKernelSURF(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3], static void openCLExecuteKernelSURF(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3],
size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels, int depth) size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels, int depth)
{ {
char * pSURF_OPTIONS = SURF_OPTIONS; char optBuf [100] = {0};
static bool OPTION_INIT = false; char * optBufPtr = optBuf;
if(!OPTION_INIT) if( !use_image2d )
{ {
if( !USE_IMAGE2d ) strcat(optBufPtr, noImage2dOption);
{ optBufPtr += strlen(noImage2dOption);
strcat(pSURF_OPTIONS, noImage2dOption);
pSURF_OPTIONS += strlen(noImage2dOption);
} }
cl_kernel kernel;
size_t wave_size = 0; kernel = openCLGetKernelFromSource(clCxt, source, kernelName, optBufPtr);
queryDeviceInfo(WAVEFRONT_SIZE, &wave_size); size_t wave_size = queryDeviceInfo<WAVEFRONT_SIZE, size_t>(kernel);
std::sprintf(pSURF_OPTIONS, "-D WAVE_SIZE=%d", static_cast<int>(wave_size)); CV_Assert(clReleaseKernel(kernel) == CL_SUCCESS);
OPTION_INIT = true; sprintf(optBufPtr, "-D WAVE_SIZE=%d", static_cast<int>(wave_size));
} openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, channels, depth, optBufPtr);
openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, channels, depth, SURF_OPTIONS);
} }
} }
} }
@ -161,22 +158,12 @@ public:
counters.setTo(Scalar::all(0)); counters.setTo(Scalar::all(0));
integral(img, surf_.sum); integral(img, surf_.sum);
if(support_image2d()) use_image2d = support_image2d();
{ if(use_image2d)
try
{ {
bindImgTex(img, imgTex); bindImgTex(img, imgTex);
bindImgTex(surf_.sum, sumTex); bindImgTex(surf_.sum, sumTex);
USE_IMAGE2d = true; finish();
}
catch (const cv::Exception& e)
{
USE_IMAGE2d = false;
if(e.code != CL_IMAGE_FORMAT_NOT_SUPPORTED && e.code != -217)
{
throw e;
}
}
} }
maskSumTex = 0; maskSumTex = 0;

View File

@ -128,11 +128,17 @@ namespace cv
enum DEVICE_INFO enum DEVICE_INFO
{ {
WAVEFRONT_SIZE, //in AMD speak WAVEFRONT_SIZE, //in AMD speak
WARP_SIZE = WAVEFRONT_SIZE, //in nvidia speak
IS_CPU_DEVICE //check if the device is CPU IS_CPU_DEVICE //check if the device is CPU
}; };
template<DEVICE_INFO _it, typename _ty>
_ty queryDeviceInfo(cl_kernel kernel = NULL);
//info should have been pre-allocated //info should have been pre-allocated
void CV_EXPORTS queryDeviceInfo(DEVICE_INFO info_type, void* info); 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 ocl

View File

@ -1578,8 +1578,7 @@ static void openCLExecuteKernel_hog(Context *clCxt , const char **source, string
size_t globalThreads[3], size_t localThreads[3], size_t globalThreads[3], size_t localThreads[3],
vector< pair<size_t, const void *> > &args) vector< pair<size_t, const void *> > &args)
{ {
size_t wave_size = 0; size_t wave_size = queryDeviceInfo<WAVEFRONT_SIZE, size_t>();
queryDeviceInfo(WAVEFRONT_SIZE, &wave_size);
if (wave_size <= 16) if (wave_size <= 16)
{ {
char build_options[64]; char build_options[64];

View File

@ -363,64 +363,43 @@ namespace cv
clFinish(Context::getContext()->impl->clCmdQueue); clFinish(Context::getContext()->impl->clCmdQueue);
} }
void queryDeviceInfo(DEVICE_INFO info_type, void* info) //template specializations of queryDeviceInfo
{ template<>
static Info::Impl* impl = Context::getContext()->impl; bool queryDeviceInfo<IS_CPU_DEVICE, bool>(cl_kernel)
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));
}
}
break;
case IS_CPU_DEVICE:
{ {
Info::Impl* impl = Context::getContext()->impl;
cl_device_type devicetype; cl_device_type devicetype;
openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum], openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum],
CL_DEVICE_TYPE, sizeof(cl_device_type), CL_DEVICE_TYPE, sizeof(cl_device_type),
&devicetype, NULL)); &devicetype, NULL));
*(bool*)info = (devicetype == CVCL_DEVICE_TYPE_CPU); return (devicetype == CVCL_DEVICE_TYPE_CPU);
} }
break;
default: template<typename _ty>
CV_Error(-1, "Invalid device info type"); static _ty queryWavesize(cl_kernel kernel)
break; {
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) void openCLReadBuffer(Context *clCxt, cl_mem dst_buffer, void *host_buffer, size_t size)

View File

@ -187,8 +187,7 @@ static void lkSparse_run(oclMat &I, oclMat &J,
args.push_back( make_pair( sizeof(cl_int), (void *)&iters )); args.push_back( make_pair( sizeof(cl_int), (void *)&iters ));
args.push_back( make_pair( sizeof(cl_char), (void *)&calcErr )); args.push_back( make_pair( sizeof(cl_char), (void *)&calcErr ));
bool is_cpu; bool is_cpu = queryDeviceInfo<IS_CPU_DEVICE, bool>();
queryDeviceInfo(IS_CPU_DEVICE, &is_cpu);
if (is_cpu) if (is_cpu)
{ {
openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), (char*)" -D CPU"); openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), (char*)" -D CPU");