From ac21cabda263776e2645796796afbd0ea84dc2e1 Mon Sep 17 00:00:00 2001 From: peng xiao Date: Tue, 14 May 2013 17:50:38 +0800 Subject: [PATCH] Copy ocl::queryDeviceInfo interface from master to 2.4. Affected functions surf.ocl, pyrlk.ocl and hog.ocl are updated with the change. --- modules/nonfree/src/surf.ocl.cpp | 49 +++++-------- modules/ocl/include/opencv2/ocl/private/util.hpp | 10 ++- modules/ocl/src/hog.cpp | 3 +- modules/ocl/src/initialization.cpp | 89 +++++++++--------------- modules/ocl/src/pyrlk.cpp | 3 +- 5 files changed, 62 insertions(+), 92 deletions(-) diff --git a/modules/nonfree/src/surf.ocl.cpp b/modules/nonfree/src/surf.ocl.cpp index acc188e..de7cac2 100644 --- a/modules/nonfree/src/surf.ocl.cpp +++ b/modules/nonfree/src/surf.ocl.cpp @@ -60,27 +60,24 @@ namespace cv 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], size_t localThreads[3], vector< pair > &args, int channels, int depth) { - char * pSURF_OPTIONS = SURF_OPTIONS; - static bool OPTION_INIT = false; - if(!OPTION_INIT) + char optBuf [100] = {0}; + char * optBufPtr = optBuf; + if( !use_image2d ) { - if( !USE_IMAGE2d ) - { - strcat(pSURF_OPTIONS, noImage2dOption); - pSURF_OPTIONS += strlen(noImage2dOption); - } - - size_t wave_size = 0; - queryDeviceInfo(WAVEFRONT_SIZE, &wave_size); - std::sprintf(pSURF_OPTIONS, "-D WAVE_SIZE=%d", static_cast(wave_size)); - OPTION_INIT = true; + strcat(optBufPtr, noImage2dOption); + optBufPtr += strlen(noImage2dOption); } - openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, channels, depth, SURF_OPTIONS); + cl_kernel kernel; + kernel = openCLGetKernelFromSource(clCxt, source, kernelName, optBufPtr); + size_t wave_size = queryDeviceInfo(kernel); + CV_Assert(clReleaseKernel(kernel) == CL_SUCCESS); + sprintf(optBufPtr, "-D WAVE_SIZE=%d", static_cast(wave_size)); + openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, channels, depth, optBufPtr); } } } @@ -161,22 +158,12 @@ public: counters.setTo(Scalar::all(0)); integral(img, surf_.sum); - if(support_image2d()) + use_image2d = support_image2d(); + if(use_image2d) { - try - { - bindImgTex(img, imgTex); - bindImgTex(surf_.sum, sumTex); - USE_IMAGE2d = true; - } - catch (const cv::Exception& e) - { - USE_IMAGE2d = false; - if(e.code != CL_IMAGE_FORMAT_NOT_SUPPORTED && e.code != -217) - { - throw e; - } - } + bindImgTex(img, imgTex); + bindImgTex(surf_.sum, sumTex); + finish(); } maskSumTex = 0; diff --git a/modules/ocl/include/opencv2/ocl/private/util.hpp b/modules/ocl/include/opencv2/ocl/private/util.hpp index 081d234..f3e582f 100644 --- a/modules/ocl/include/opencv2/ocl/private/util.hpp +++ b/modules/ocl/include/opencv2/ocl/private/util.hpp @@ -128,11 +128,17 @@ namespace cv enum DEVICE_INFO { WAVEFRONT_SIZE, //in AMD speak - WARP_SIZE = WAVEFRONT_SIZE, //in nvidia speak IS_CPU_DEVICE //check if the device is CPU }; + template + _ty queryDeviceInfo(cl_kernel kernel = NULL); //info should have been pre-allocated - void CV_EXPORTS queryDeviceInfo(DEVICE_INFO info_type, void* info); + template<> + int CV_EXPORTS queryDeviceInfo(cl_kernel kernel); + template<> + size_t CV_EXPORTS queryDeviceInfo(cl_kernel kernel); + template<> + bool CV_EXPORTS queryDeviceInfo(cl_kernel kernel); }//namespace ocl diff --git a/modules/ocl/src/hog.cpp b/modules/ocl/src/hog.cpp index 7a13324..9c8f315 100644 --- a/modules/ocl/src/hog.cpp +++ b/modules/ocl/src/hog.cpp @@ -1578,8 +1578,7 @@ static void openCLExecuteKernel_hog(Context *clCxt , const char **source, string size_t globalThreads[3], size_t localThreads[3], vector< pair > &args) { - size_t wave_size = 0; - queryDeviceInfo(WAVEFRONT_SIZE, &wave_size); + size_t wave_size = queryDeviceInfo(); if (wave_size <= 16) { char build_options[64]; diff --git a/modules/ocl/src/initialization.cpp b/modules/ocl/src/initialization.cpp index 799c49c..fd462ad 100644 --- a/modules/ocl/src/initialization.cpp +++ b/modules/ocl/src/initialization.cpp @@ -363,64 +363,43 @@ namespace cv clFinish(Context::getContext()->impl->clCmdQueue); } - void queryDeviceInfo(DEVICE_INFO info_type, void* info) + //template specializations of queryDeviceInfo + template<> + bool queryDeviceInfo(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 + static _ty queryWavesize(cl_kernel kernel) + { + size_t info = 0; + Info::Impl* impl = Context::getContext()->impl; + bool is_cpu = queryDeviceInfo(); + 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(cl_kernel kernel) + { + return queryWavesize(kernel); + } + template<> + int queryDeviceInfo(cl_kernel kernel) + { + return queryWavesize(kernel); } void openCLReadBuffer(Context *clCxt, cl_mem dst_buffer, void *host_buffer, size_t size) diff --git a/modules/ocl/src/pyrlk.cpp b/modules/ocl/src/pyrlk.cpp index 4a6ce1c..6de4f97 100644 --- a/modules/ocl/src/pyrlk.cpp +++ b/modules/ocl/src/pyrlk.cpp @@ -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_char), (void *)&calcErr )); - bool is_cpu; - queryDeviceInfo(IS_CPU_DEVICE, &is_cpu); + bool is_cpu = queryDeviceInfo(); if (is_cpu) { openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), (char*)" -D CPU"); -- 2.7.4