其他分享
首页 > 其他分享> > OpenCL For Opencv

OpenCL For Opencv

作者:互联网

一、Opencv-OCL编程基础

1. Opencv OCL基本编程API

2. 图像处理Kernel实现及CU单元配置

3. Demo实验

        我目前编译使用的opencv版本是opencv4.2【如果版本不同请下载不同版本下的Demo程序】,使用如下官方提供的Opencv-OCL代码,如果只是单纯的测试运行此官方提供的代码不需要有特定的加速设备,直接使用多核心CPU-PC平台即可,因为OpenCL本身就支持了CPU加速,具体代码如下:

  1 // This file is part of OpenCV project.
  2 // It is subject to the license terms in the LICENSE file found in the top-level directory
  3 // of this distribution and at http://opencv.org/license.html
  4 
  5 #include "opencv2/core.hpp"
  6 #include "opencv2/core/ocl.hpp"
  7 #include "opencv2/highgui.hpp"
  8 #include "opencv2/imgcodecs.hpp"
  9 #include "opencv2/imgproc.hpp"
 10 
 11 #include <iostream>
 12 
 13 using namespace std;
 14 using namespace cv;
 15 
 16 static const char* opencl_kernel_src =
 17 "__kernel void magnutude_filter_8u(\n"
 18 "       __global const uchar* src, int src_step, int src_offset,\n"
 19 "       __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n"
 20 "       float scale)\n"
 21 "{\n"
 22 "   int x = get_global_id(0);\n"
 23 "   int y = get_global_id(1);\n"
 24 "   if (x < dst_cols && y < dst_rows)\n"
 25 "   {\n"
 26 "       int dst_idx = y * dst_step + x + dst_offset;\n"
 27 "       if (x > 0 && x < dst_cols - 1 && y > 0 && y < dst_rows - 2)\n"
 28 "       {\n"
 29 "           int src_idx = y * src_step + x + src_offset;\n"
 30 "           int dx = (int)src[src_idx]*2 - src[src_idx - 1]          - src[src_idx + 1];\n"
 31 "           int dy = (int)src[src_idx]*2 - src[src_idx - 1*src_step] - src[src_idx + 1*src_step];\n"
 32 "           dst[dst_idx] = convert_uchar_sat(sqrt((float)(dx*dx + dy*dy)) * scale);\n"
 33 "       }\n"
 34 "       else\n"
 35 "       {\n"
 36 "           dst[dst_idx] = 0;\n"
 37 "       }\n"
 38 "   }\n"
 39 "}\n";
 40 
 41 int main(int argc, char** argv)
 42 {
 43     const char* keys =
 44         "{ i input    | | specify input image }"
 45         "{ h help     | | print help message }";
 46 
 47     cv::CommandLineParser args(argc, argv, keys);
 48     if (args.has("help"))
 49     {
 50         cout << "Usage : " << argv[0] << " [options]" << endl;
 51         cout << "Available options:" << endl;
 52         args.printMessage();
 53         return EXIT_SUCCESS;
 54     }
 55 
 56     cv::ocl::Context ctx = cv::ocl::Context::getDefault();
 57     if (!ctx.ptr())
 58     {
 59         cerr << "OpenCL is not available" << endl;
 60         return 1;
 61     }
 62     cv::ocl::Device device = cv::ocl::Device::getDefault();
 63     if (!device.compilerAvailable())
 64     {
 65         cerr << "OpenCL compiler is not available" << endl;
 66         return 1;
 67     }
 68 
 69 
 70     UMat src;
 71     {
 72         string image_file = args.get<string>("i");
 73         if (!image_file.empty())
 74         {
 75             Mat image = imread(samples::findFile(image_file));
 76             if (image.empty())
 77             {
 78                 cout << "error read image: " << image_file << endl;
 79                 return 1;
 80             }
 81             cvtColor(image, src, COLOR_BGR2GRAY);
 82         }
 83         else
 84         {
 85             Mat frame(cv::Size(640, 480), CV_8U, Scalar::all(128));
 86             Point p(frame.cols / 2, frame.rows / 2);
 87             line(frame, Point(0, frame.rows / 2), Point(frame.cols, frame.rows / 2), 1);
 88             circle(frame, p, 200, Scalar(32, 32, 32), 8, LINE_AA);
 89             string str = "OpenCL";
 90             int baseLine = 0;
 91             Size box = getTextSize(str, FONT_HERSHEY_COMPLEX, 2, 5, &baseLine);
 92             putText(frame, str, Point((frame.cols - box.width) / 2, (frame.rows - box.height) / 2 + baseLine),
 93                     FONT_HERSHEY_COMPLEX, 2, Scalar(255, 255, 255), 5, LINE_AA);
 94             frame.copyTo(src);
 95         }
 96     }
 97 
 98 
 99     cv::String module_name; // empty to disable OpenCL cache
100 
101     {
102         cout << "OpenCL program source: " << endl;
103         cout << "======================================================================================================" << endl;
104         cout << opencl_kernel_src << endl;
105         cout << "======================================================================================================" << endl;
106         //! [Define OpenCL program source]
107         cv::ocl::ProgramSource source(module_name, "simple", opencl_kernel_src, "");
108         //! [Define OpenCL program source]
109 
110         //! [Compile/build OpenCL for current OpenCL device]
111         cv::String errmsg;
112         cv::ocl::Program program(source, "", errmsg);
113         if (program.ptr() == NULL)
114         {
115             cerr << "Can't compile OpenCL program:" << endl << errmsg << endl;
116             return 1;
117         }
118         //! [Compile/build OpenCL for current OpenCL device]
119 
120         if (!errmsg.empty())
121         {
122             cout << "OpenCL program build log:" << endl << errmsg << endl;
123         }
124 
125         //! [Get OpenCL kernel by name]
126         cv::ocl::Kernel k("magnutude_filter_8u", program);
127         if (k.empty())
128         {
129             cerr << "Can't get OpenCL kernel" << endl;
130             return 1;
131         }
132         //! [Get OpenCL kernel by name]
133 
134         UMat result(src.size(), CV_8UC1);
135 
136         //! [Define kernel parameters and run]
137         size_t globalSize[2] = {(size_t)src.cols, (size_t)src.rows};
138         size_t localSize[2] = {8, 8};
139         bool executionResult = k
140             .args(
141                 cv::ocl::KernelArg::ReadOnlyNoSize(src), // size is not used (similar to 'dst' size)
142                 cv::ocl::KernelArg::WriteOnly(result),
143                 (float)2.0
144             )
145             .run(2, globalSize, localSize, true);
146         if (!executionResult)
147         {
148             cerr << "OpenCL kernel launch failed" << endl;
149             return 1;
150         }
151         //! [Define kernel parameters and run]
152 
153         imshow("Source", src);
154         imshow("Result", result);
155 
156         for (;;)
157         {
158             int key = waitKey();
159             if (key == 27/*ESC*/ || key == 'q' || key == 'Q')
160                 break;
161         }
162     }
163     return 0;
164 }

使用mingw编译上述程序并运行结果如下(根据运行结果,说明对图像进行了边缘提取的功能):

当然如果你想处理其他的图像,也可以使用在CMD窗口当中调用编译完成的.exe文件加上文件的绝对路径。

终端Terminal当中显示了Kernel的基本内容,如下所示:

        从Kernel的结构可以分析出来,kernel中实现的是求解了图像自身的x方向以及y方向的梯度,并求出了每一点的梯度方向,实际上比较类似与Canny边缘算子的检测算法,kernel使用了二维的方式处理图片上每一个点,在kernel核当中,使用了if判断是否存在指针越界的情况,具体的基础实现相关内容请移步OpenCL基础入门

二、Demo代码变形

实现图像3x3均值滤波

三、嵌入式平台移植与编译(TI AM57x 系列)

Opencv3.1.0版本master下,opencv-ocl代码官方Demo

   1 /*
   2 // The example of interoperability between OpenCL and OpenCV.
   3 // This will loop through frames of video either from input media file
   4 // or camera device and do processing of these data in OpenCL and then
   5 // in OpenCV. In OpenCL it does inversion of pixels in left half of frame and
   6 // in OpenCV it does bluring in the right half of frame.
   7 */
   8 #include <cstdio>
   9 #include <cstdlib>
  10 #include <iostream>
  11 #include <fstream>
  12 #include <string>
  13 #include <sstream>
  14 #include <iomanip>
  15 #include <stdexcept>
  16 
  17 #define CL_USE_DEPRECATED_OPENCL_2_0_APIS // eliminate build warning
  18 
  19 #if __APPLE__
  20 #include <OpenCL/cl.h>
  21 #else
  22 #include <CL/cl.h>
  23 #endif
  24 
  25 #include <opencv2/core/ocl.hpp>
  26 #include <opencv2/core/utility.hpp>
  27 #include <opencv2/video.hpp>
  28 #include <opencv2/highgui.hpp>
  29 #include <opencv2/imgproc.hpp>
  30 
  31 
  32 using namespace std;
  33 using namespace cv;
  34 
  35 namespace opencl {
  36 
  37 class PlatformInfo
  38 {
  39 public:
  40     PlatformInfo()
  41     {}
  42 
  43     ~PlatformInfo()
  44     {}
  45 
  46     cl_int QueryInfo(cl_platform_id id)
  47     {
  48         query_param(id, CL_PLATFORM_PROFILE, m_profile);
  49         query_param(id, CL_PLATFORM_VERSION, m_version);
  50         query_param(id, CL_PLATFORM_NAME, m_name);
  51         query_param(id, CL_PLATFORM_VENDOR, m_vendor);
  52         query_param(id, CL_PLATFORM_EXTENSIONS, m_extensions);
  53         return CL_SUCCESS;
  54     }
  55 
  56     std::string Profile()    { return m_profile; }
  57     std::string Version()    { return m_version; }
  58     std::string Name()       { return m_name; }
  59     std::string Vendor()     { return m_vendor; }
  60     std::string Extensions() { return m_extensions; }
  61 
  62 private:
  63     cl_int query_param(cl_platform_id id, cl_platform_info param, std::string& paramStr)
  64     {
  65         cl_int res;
  66 
  67         size_t psize;
  68         cv::AutoBuffer<char> buf;
  69 
  70         res = clGetPlatformInfo(id, param, 0, 0, &psize);
  71         if (CL_SUCCESS != res)
  72             throw std::runtime_error(std::string("clGetPlatformInfo failed"));
  73 
  74         buf.resize(psize);
  75         res = clGetPlatformInfo(id, param, psize, buf, 0);
  76         if (CL_SUCCESS != res)
  77             throw std::runtime_error(std::string("clGetPlatformInfo failed"));
  78 
  79         // just in case, ensure trailing zero for ASCIIZ string
  80         buf[psize] = 0;
  81 
  82         paramStr = buf;
  83 
  84         return CL_SUCCESS;
  85     }
  86 
  87 private:
  88     std::string m_profile;
  89     std::string m_version;
  90     std::string m_name;
  91     std::string m_vendor;
  92     std::string m_extensions;
  93 };
  94 
  95 
  96 class DeviceInfo
  97 {
  98 public:
  99     DeviceInfo()
 100     {}
 101 
 102     ~DeviceInfo()
 103     {}
 104 
 105     cl_int QueryInfo(cl_device_id id)
 106     {
 107         query_param(id, CL_DEVICE_TYPE, m_type);
 108         query_param(id, CL_DEVICE_VENDOR_ID, m_vendor_id);
 109         query_param(id, CL_DEVICE_MAX_COMPUTE_UNITS, m_max_compute_units);
 110         query_param(id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, m_max_work_item_dimensions);
 111         query_param(id, CL_DEVICE_MAX_WORK_ITEM_SIZES, m_max_work_item_sizes);
 112         query_param(id, CL_DEVICE_MAX_WORK_GROUP_SIZE, m_max_work_group_size);
 113         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, m_preferred_vector_width_char);
 114         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, m_preferred_vector_width_short);
 115         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, m_preferred_vector_width_int);
 116         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, m_preferred_vector_width_long);
 117         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, m_preferred_vector_width_float);
 118         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, m_preferred_vector_width_double);
 119 #if defined(CL_VERSION_1_1)
 120         query_param(id, CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF, m_preferred_vector_width_half);
 121         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR, m_native_vector_width_char);
 122         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT, m_native_vector_width_short);
 123         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, m_native_vector_width_int);
 124         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, m_native_vector_width_long);
 125         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT, m_native_vector_width_float);
 126         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE, m_native_vector_width_double);
 127         query_param(id, CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF, m_native_vector_width_half);
 128 #endif
 129         query_param(id, CL_DEVICE_MAX_CLOCK_FREQUENCY, m_max_clock_frequency);
 130         query_param(id, CL_DEVICE_ADDRESS_BITS, m_address_bits);
 131         query_param(id, CL_DEVICE_MAX_MEM_ALLOC_SIZE, m_max_mem_alloc_size);
 132         query_param(id, CL_DEVICE_IMAGE_SUPPORT, m_image_support);
 133         query_param(id, CL_DEVICE_MAX_READ_IMAGE_ARGS, m_max_read_image_args);
 134         query_param(id, CL_DEVICE_MAX_WRITE_IMAGE_ARGS, m_max_write_image_args);
 135 #if defined(CL_VERSION_2_0)
 136         query_param(id, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS, m_max_read_write_image_args);
 137 #endif
 138         query_param(id, CL_DEVICE_IMAGE2D_MAX_WIDTH, m_image2d_max_width);
 139         query_param(id, CL_DEVICE_IMAGE2D_MAX_HEIGHT, m_image2d_max_height);
 140         query_param(id, CL_DEVICE_IMAGE3D_MAX_WIDTH, m_image3d_max_width);
 141         query_param(id, CL_DEVICE_IMAGE3D_MAX_HEIGHT, m_image3d_max_height);
 142         query_param(id, CL_DEVICE_IMAGE3D_MAX_DEPTH, m_image3d_max_depth);
 143 #if defined(CL_VERSION_1_2)
 144         query_param(id, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, m_image_max_buffer_size);
 145         query_param(id, CL_DEVICE_IMAGE_MAX_ARRAY_SIZE, m_image_max_array_size);
 146 #endif
 147         query_param(id, CL_DEVICE_MAX_SAMPLERS, m_max_samplers);
 148 #if defined(CL_VERSION_1_2)
 149         query_param(id, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, m_image_pitch_alignment);
 150         query_param(id, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, m_image_base_address_alignment);
 151 #endif
 152 #if defined(CL_VERSION_2_0)
 153         query_param(id, CL_DEVICE_MAX_PIPE_ARGS, m_max_pipe_args);
 154         query_param(id, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, m_pipe_max_active_reservations);
 155         query_param(id, CL_DEVICE_PIPE_MAX_PACKET_SIZE, m_pipe_max_packet_size);
 156 #endif
 157         query_param(id, CL_DEVICE_MAX_PARAMETER_SIZE, m_max_parameter_size);
 158         query_param(id, CL_DEVICE_MEM_BASE_ADDR_ALIGN, m_mem_base_addr_align);
 159         query_param(id, CL_DEVICE_SINGLE_FP_CONFIG, m_single_fp_config);
 160 #if defined(CL_VERSION_1_2)
 161         query_param(id, CL_DEVICE_DOUBLE_FP_CONFIG, m_double_fp_config);
 162 #endif
 163         query_param(id, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, m_global_mem_cache_type);
 164         query_param(id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, m_global_mem_cacheline_size);
 165         query_param(id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, m_global_mem_cache_size);
 166         query_param(id, CL_DEVICE_GLOBAL_MEM_SIZE, m_global_mem_size);
 167         query_param(id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, m_max_constant_buffer_size);
 168         query_param(id, CL_DEVICE_MAX_CONSTANT_ARGS, m_max_constant_args);
 169 #if defined(CL_VERSION_2_0)
 170         query_param(id, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, m_max_global_variable_size);
 171         query_param(id, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, m_global_variable_preferred_total_size);
 172 #endif
 173         query_param(id, CL_DEVICE_LOCAL_MEM_TYPE, m_local_mem_type);
 174         query_param(id, CL_DEVICE_LOCAL_MEM_SIZE, m_local_mem_size);
 175         query_param(id, CL_DEVICE_ERROR_CORRECTION_SUPPORT, m_error_correction_support);
 176 #if defined(CL_VERSION_1_1)
 177         query_param(id, CL_DEVICE_HOST_UNIFIED_MEMORY, m_host_unified_memory);
 178 #endif
 179         query_param(id, CL_DEVICE_PROFILING_TIMER_RESOLUTION, m_profiling_timer_resolution);
 180         query_param(id, CL_DEVICE_ENDIAN_LITTLE, m_endian_little);
 181         query_param(id, CL_DEVICE_AVAILABLE, m_available);
 182         query_param(id, CL_DEVICE_COMPILER_AVAILABLE, m_compiler_available);
 183 #if defined(CL_VERSION_1_2)
 184         query_param(id, CL_DEVICE_LINKER_AVAILABLE, m_linker_available);
 185 #endif
 186         query_param(id, CL_DEVICE_EXECUTION_CAPABILITIES, m_execution_capabilities);
 187         query_param(id, CL_DEVICE_QUEUE_PROPERTIES, m_queue_properties);
 188 #if defined(CL_VERSION_2_0)
 189         query_param(id, CL_DEVICE_QUEUE_ON_HOST_PROPERTIES, m_queue_on_host_properties);
 190         query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES, m_queue_on_device_properties);
 191         query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE, m_queue_on_device_preferred_size);
 192         query_param(id, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, m_queue_on_device_max_size);
 193         query_param(id, CL_DEVICE_MAX_ON_DEVICE_QUEUES, m_max_on_device_queues);
 194         query_param(id, CL_DEVICE_MAX_ON_DEVICE_EVENTS, m_max_on_device_events);
 195 #endif
 196 #if defined(CL_VERSION_1_2)
 197         query_param(id, CL_DEVICE_BUILT_IN_KERNELS, m_built_in_kernels);
 198 #endif
 199         query_param(id, CL_DEVICE_PLATFORM, m_platform);
 200         query_param(id, CL_DEVICE_NAME, m_name);
 201         query_param(id, CL_DEVICE_VENDOR, m_vendor);
 202         query_param(id, CL_DRIVER_VERSION, m_driver_version);
 203         query_param(id, CL_DEVICE_PROFILE, m_profile);
 204         query_param(id, CL_DEVICE_VERSION, m_version);
 205 #if defined(CL_VERSION_1_1)
 206         query_param(id, CL_DEVICE_OPENCL_C_VERSION, m_opencl_c_version);
 207 #endif
 208         query_param(id, CL_DEVICE_EXTENSIONS, m_extensions);
 209 #if defined(CL_VERSION_1_2)
 210         query_param(id, CL_DEVICE_PRINTF_BUFFER_SIZE, m_printf_buffer_size);
 211         query_param(id, CL_DEVICE_PREFERRED_INTEROP_USER_SYNC, m_preferred_interop_user_sync);
 212         query_param(id, CL_DEVICE_PARENT_DEVICE, m_parent_device);
 213         query_param(id, CL_DEVICE_PARTITION_MAX_SUB_DEVICES, m_partition_max_sub_devices);
 214         query_param(id, CL_DEVICE_PARTITION_PROPERTIES, m_partition_properties);
 215         query_param(id, CL_DEVICE_PARTITION_AFFINITY_DOMAIN, m_partition_affinity_domain);
 216         query_param(id, CL_DEVICE_PARTITION_TYPE, m_partition_type);
 217         query_param(id, CL_DEVICE_REFERENCE_COUNT, m_reference_count);
 218 #endif
 219         return CL_SUCCESS;
 220     }
 221 
 222     std::string Name() { return m_name; }
 223 
 224 private:
 225     template<typename T>
 226     cl_int query_param(cl_device_id id, cl_device_info param, T& value)
 227     {
 228         cl_int res;
 229         size_t size = 0;
 230 
 231         res = clGetDeviceInfo(id, param, 0, 0, &size);
 232         if (CL_SUCCESS != res && size != 0)
 233             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
 234 
 235         if (0 == size)
 236             return CL_SUCCESS;
 237 
 238         if (sizeof(T) != size)
 239             throw std::runtime_error(std::string("clGetDeviceInfo: param size mismatch"));
 240 
 241         res = clGetDeviceInfo(id, param, size, &value, 0);
 242         if (CL_SUCCESS != res)
 243             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
 244 
 245         return CL_SUCCESS;
 246     }
 247 
 248     template<typename T>
 249     cl_int query_param(cl_device_id id, cl_device_info param, std::vector<T>& value)
 250     {
 251         cl_int res;
 252         size_t size;
 253 
 254         res = clGetDeviceInfo(id, param, 0, 0, &size);
 255         if (CL_SUCCESS != res)
 256             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
 257 
 258         if (0 == size)
 259             return CL_SUCCESS;
 260 
 261         value.resize(size / sizeof(T));
 262 
 263         res = clGetDeviceInfo(id, param, size, &value[0], 0);
 264         if (CL_SUCCESS != res)
 265             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
 266 
 267         return CL_SUCCESS;
 268     }
 269 
 270     cl_int query_param(cl_device_id id, cl_device_info param, std::string& value)
 271     {
 272         cl_int res;
 273         size_t size;
 274 
 275         res = clGetDeviceInfo(id, param, 0, 0, &size);
 276         if (CL_SUCCESS != res)
 277             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
 278 
 279         value.resize(size + 1);
 280 
 281         res = clGetDeviceInfo(id, param, size, &value[0], 0);
 282         if (CL_SUCCESS != res)
 283             throw std::runtime_error(std::string("clGetDeviceInfo failed"));
 284 
 285         // just in case, ensure trailing zero for ASCIIZ string
 286         value[size] = 0;
 287 
 288         return CL_SUCCESS;
 289     }
 290 
 291 private:
 292     cl_device_type                            m_type;
 293     cl_uint                                   m_vendor_id;
 294     cl_uint                                   m_max_compute_units;
 295     cl_uint                                   m_max_work_item_dimensions;
 296     std::vector<size_t>                       m_max_work_item_sizes;
 297     size_t                                    m_max_work_group_size;
 298     cl_uint                                   m_preferred_vector_width_char;
 299     cl_uint                                   m_preferred_vector_width_short;
 300     cl_uint                                   m_preferred_vector_width_int;
 301     cl_uint                                   m_preferred_vector_width_long;
 302     cl_uint                                   m_preferred_vector_width_float;
 303     cl_uint                                   m_preferred_vector_width_double;
 304 #if defined(CL_VERSION_1_1)
 305     cl_uint                                   m_preferred_vector_width_half;
 306     cl_uint                                   m_native_vector_width_char;
 307     cl_uint                                   m_native_vector_width_short;
 308     cl_uint                                   m_native_vector_width_int;
 309     cl_uint                                   m_native_vector_width_long;
 310     cl_uint                                   m_native_vector_width_float;
 311     cl_uint                                   m_native_vector_width_double;
 312     cl_uint                                   m_native_vector_width_half;
 313 #endif
 314     cl_uint                                   m_max_clock_frequency;
 315     cl_uint                                   m_address_bits;
 316     cl_ulong                                  m_max_mem_alloc_size;
 317     cl_bool                                   m_image_support;
 318     cl_uint                                   m_max_read_image_args;
 319     cl_uint                                   m_max_write_image_args;
 320 #if defined(CL_VERSION_2_0)
 321     cl_uint                                   m_max_read_write_image_args;
 322 #endif
 323     size_t                                    m_image2d_max_width;
 324     size_t                                    m_image2d_max_height;
 325     size_t                                    m_image3d_max_width;
 326     size_t                                    m_image3d_max_height;
 327     size_t                                    m_image3d_max_depth;
 328 #if defined(CL_VERSION_1_2)
 329     size_t                                    m_image_max_buffer_size;
 330     size_t                                    m_image_max_array_size;
 331 #endif
 332     cl_uint                                   m_max_samplers;
 333 #if defined(CL_VERSION_1_2)
 334     cl_uint                                   m_image_pitch_alignment;
 335     cl_uint                                   m_image_base_address_alignment;
 336 #endif
 337 #if defined(CL_VERSION_2_0)
 338     cl_uint                                   m_max_pipe_args;
 339     cl_uint                                   m_pipe_max_active_reservations;
 340     cl_uint                                   m_pipe_max_packet_size;
 341 #endif
 342     size_t                                    m_max_parameter_size;
 343     cl_uint                                   m_mem_base_addr_align;
 344     cl_device_fp_config                       m_single_fp_config;
 345 #if defined(CL_VERSION_1_2)
 346     cl_device_fp_config                       m_double_fp_config;
 347 #endif
 348     cl_device_mem_cache_type                  m_global_mem_cache_type;
 349     cl_uint                                   m_global_mem_cacheline_size;
 350     cl_ulong                                  m_global_mem_cache_size;
 351     cl_ulong                                  m_global_mem_size;
 352     cl_ulong                                  m_max_constant_buffer_size;
 353     cl_uint                                   m_max_constant_args;
 354 #if defined(CL_VERSION_2_0)
 355     size_t                                    m_max_global_variable_size;
 356     size_t                                    m_global_variable_preferred_total_size;
 357 #endif
 358     cl_device_local_mem_type                  m_local_mem_type;
 359     cl_ulong                                  m_local_mem_size;
 360     cl_bool                                   m_error_correction_support;
 361 #if defined(CL_VERSION_1_1)
 362     cl_bool                                   m_host_unified_memory;
 363 #endif
 364     size_t                                    m_profiling_timer_resolution;
 365     cl_bool                                   m_endian_little;
 366     cl_bool                                   m_available;
 367     cl_bool                                   m_compiler_available;
 368 #if defined(CL_VERSION_1_2)
 369     cl_bool                                   m_linker_available;
 370 #endif
 371     cl_device_exec_capabilities               m_execution_capabilities;
 372     cl_command_queue_properties               m_queue_properties;
 373 #if defined(CL_VERSION_2_0)
 374     cl_command_queue_properties               m_queue_on_host_properties;
 375     cl_command_queue_properties               m_queue_on_device_properties;
 376     cl_uint                                   m_queue_on_device_preferred_size;
 377     cl_uint                                   m_queue_on_device_max_size;
 378     cl_uint                                   m_max_on_device_queues;
 379     cl_uint                                   m_max_on_device_events;
 380 #endif
 381 #if defined(CL_VERSION_1_2)
 382     std::string                               m_built_in_kernels;
 383 #endif
 384     cl_platform_id                            m_platform;
 385     std::string                               m_name;
 386     std::string                               m_vendor;
 387     std::string                               m_driver_version;
 388     std::string                               m_profile;
 389     std::string                               m_version;
 390 #if defined(CL_VERSION_1_1)
 391     std::string                               m_opencl_c_version;
 392 #endif
 393     std::string                               m_extensions;
 394 #if defined(CL_VERSION_1_2)
 395     size_t                                    m_printf_buffer_size;
 396     cl_bool                                   m_preferred_interop_user_sync;
 397     cl_device_id                              m_parent_device;
 398     cl_uint                                   m_partition_max_sub_devices;
 399     std::vector<cl_device_partition_property> m_partition_properties;
 400     cl_device_affinity_domain                 m_partition_affinity_domain;
 401     std::vector<cl_device_partition_property> m_partition_type;
 402     cl_uint                                   m_reference_count;
 403 #endif
 404 };
 405 
 406 } // namespace opencl
 407 
 408 
 409 class App
 410 {
 411 public:
 412     App(CommandLineParser& cmd);
 413     ~App();
 414 
 415     int initOpenCL();
 416     int initVideoSource();
 417 
 418     int process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* cl_buffer);
 419     int process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u);
 420     int process_cl_image_with_opencv(cl_mem image, cv::UMat& u);
 421 
 422     int run();
 423 
 424     bool isRunning() { return m_running; }
 425     bool doProcess() { return m_process; }
 426     bool useBuffer() { return m_use_buffer; }
 427 
 428     void setRunning(bool running)      { m_running = running; }
 429     void setDoProcess(bool process)    { m_process = process; }
 430     void setUseBuffer(bool use_buffer) { m_use_buffer = use_buffer; }
 431 
 432 protected:
 433     bool nextFrame(cv::Mat& frame) { return m_cap.read(frame); }
 434     void handleKey(char key);
 435     void timerStart();
 436     void timerEnd();
 437     std::string timeStr() const;
 438     std::string message() const;
 439 
 440 private:
 441     bool                        m_running;
 442     bool                        m_process;
 443     bool                        m_use_buffer;
 444 
 445     int64                       m_t0;
 446     int64                       m_t1;
 447     float                       m_time;
 448     float                       m_frequency;
 449 
 450     string                      m_file_name;
 451     int                         m_camera_id;
 452     cv::VideoCapture            m_cap;
 453     cv::Mat                     m_frame;
 454     cv::Mat                     m_frameGray;
 455 
 456     opencl::PlatformInfo        m_platformInfo;
 457     opencl::DeviceInfo          m_deviceInfo;
 458     std::vector<cl_platform_id> m_platform_ids;
 459     cl_context                  m_context;
 460     cl_device_id                m_device_id;
 461     cl_command_queue            m_queue;
 462     cl_program                  m_program;
 463     cl_kernel                   m_kernelBuf;
 464     cl_kernel                   m_kernelImg;
 465     cl_mem                      m_img_src; // used as src in case processing of cl image
 466     cl_mem                      m_mem_obj;
 467     cl_event                    m_event;
 468 };
 469 
 470 
 471 App::App(CommandLineParser& cmd)
 472 {
 473     cout << "\nPress ESC to exit\n" << endl;
 474     cout << "\n      'p' to toggle ON/OFF processing\n" << endl;
 475     cout << "\n       SPACE to switch between OpenCL buffer/image\n" << endl;
 476 
 477     m_camera_id  = cmd.get<int>("camera");
 478     m_file_name  = cmd.get<string>("video");
 479 
 480     m_running    = false;
 481     m_process    = false;
 482     m_use_buffer = false;
 483 
 484     m_t0         = 0;
 485     m_t1         = 0;
 486     m_time       = 0.0;
 487     m_frequency  = (float)cv::getTickFrequency();
 488 
 489     m_context    = 0;
 490     m_device_id  = 0;
 491     m_queue      = 0;
 492     m_program    = 0;
 493     m_kernelBuf  = 0;
 494     m_kernelImg  = 0;
 495     m_img_src    = 0;
 496     m_mem_obj    = 0;
 497     m_event      = 0;
 498 } // ctor
 499 
 500 
 501 App::~App()
 502 {
 503     if (m_queue)
 504     {
 505         clFinish(m_queue);
 506         clReleaseCommandQueue(m_queue);
 507         m_queue = 0;
 508     }
 509 
 510     if (m_program)
 511     {
 512         clReleaseProgram(m_program);
 513         m_program = 0;
 514     }
 515 
 516     if (m_img_src)
 517     {
 518         clReleaseMemObject(m_img_src);
 519         m_img_src = 0;
 520     }
 521 
 522     if (m_mem_obj)
 523     {
 524         clReleaseMemObject(m_mem_obj);
 525         m_mem_obj = 0;
 526     }
 527 
 528     if (m_event)
 529     {
 530         clReleaseEvent(m_event);
 531     }
 532 
 533     if (m_kernelBuf)
 534     {
 535         clReleaseKernel(m_kernelBuf);
 536         m_kernelBuf = 0;
 537     }
 538 
 539     if (m_kernelImg)
 540     {
 541         clReleaseKernel(m_kernelImg);
 542         m_kernelImg = 0;
 543     }
 544 
 545     if (m_device_id)
 546     {
 547         clReleaseDevice(m_device_id);
 548         m_device_id = 0;
 549     }
 550 
 551     if (m_context)
 552     {
 553         clReleaseContext(m_context);
 554         m_context = 0;
 555     }
 556 } // dtor
 557 
 558 
 559 int App::initOpenCL()
 560 {
 561     cl_int res = CL_SUCCESS;
 562     cl_uint num_entries = 0;
 563 
 564     res = clGetPlatformIDs(0, 0, &num_entries);
 565     if (CL_SUCCESS != res)
 566         return -1;
 567 
 568     m_platform_ids.resize(num_entries);
 569 
 570     res = clGetPlatformIDs(num_entries, &m_platform_ids[0], 0);
 571     if (CL_SUCCESS != res)
 572         return -1;
 573 
 574     unsigned int i;
 575 
 576     // create context from first platform with GPU device
 577     for (i = 0; i < m_platform_ids.size(); i++)
 578     {
 579         cl_context_properties props[] =
 580         {
 581             CL_CONTEXT_PLATFORM,
 582             (cl_context_properties)(m_platform_ids[i]),
 583             0
 584         };
 585 
 586         m_context = clCreateContextFromType(props, CL_DEVICE_TYPE_GPU, 0, 0, &res);
 587         if (0 == m_context || CL_SUCCESS != res)
 588             continue;
 589 
 590         res = clGetContextInfo(m_context, CL_CONTEXT_DEVICES, sizeof(cl_device_id), &m_device_id, 0);
 591         if (CL_SUCCESS != res)
 592             return -1;
 593 
 594         m_queue = clCreateCommandQueue(m_context, m_device_id, 0, &res);
 595         if (0 == m_queue || CL_SUCCESS != res)
 596             return -1;
 597 
 598         const char* kernelSrc =
 599             "__kernel "
 600             "void bitwise_inv_buf_8uC1("
 601             "    __global unsigned char* pSrcDst,"
 602             "             int            srcDstStep,"
 603             "             int            rows,"
 604             "             int            cols)"
 605             "{"
 606             "    int x = get_global_id(0);"
 607             "    int y = get_global_id(1);"
 608             "    int idx = mad24(y, srcDstStep, x);"
 609             "    pSrcDst[idx] = ~pSrcDst[idx];"
 610             "}"
 611             "__kernel "
 612             "void bitwise_inv_img_8uC1("
 613             "    read_only  image2d_t srcImg,"
 614             "    write_only image2d_t dstImg)"
 615             "{"
 616             "    int x = get_global_id(0);"
 617             "    int y = get_global_id(1);"
 618             "    int2 coord = (int2)(x, y);"
 619             "    uint4 val = read_imageui(srcImg, coord);"
 620             "    val.x = (~val.x) & 0x000000FF;"
 621             "    write_imageui(dstImg, coord, val);"
 622             "}";
 623         size_t len = strlen(kernelSrc);
 624         m_program = clCreateProgramWithSource(m_context, 1, &kernelSrc, &len, &res);
 625         if (0 == m_program || CL_SUCCESS != res)
 626             return -1;
 627 
 628         res = clBuildProgram(m_program, 1, &m_device_id, 0, 0, 0);
 629         if (CL_SUCCESS != res)
 630             return -1;
 631 
 632         m_kernelBuf = clCreateKernel(m_program, "bitwise_inv_buf_8uC1", &res);
 633         if (0 == m_kernelBuf || CL_SUCCESS != res)
 634             return -1;
 635 
 636         m_kernelImg = clCreateKernel(m_program, "bitwise_inv_img_8uC1", &res);
 637         if (0 == m_kernelImg || CL_SUCCESS != res)
 638             return -1;
 639 
 640         m_platformInfo.QueryInfo(m_platform_ids[i]);
 641         m_deviceInfo.QueryInfo(m_device_id);
 642 
 643         // attach OpenCL context to OpenCV
 644         cv::ocl::attachContext(m_platformInfo.Name(), m_platform_ids[i], m_context, m_device_id);
 645 
 646         break;
 647     }
 648 
 649     return m_context != 0 ? CL_SUCCESS : -1;
 650 } // initOpenCL()
 651 
 652 
 653 int App::initVideoSource()
 654 {
 655     try
 656     {
 657         if (!m_file_name.empty() && m_camera_id == -1)
 658         {
 659             m_cap.open(m_file_name.c_str());
 660             if (!m_cap.isOpened())
 661                 throw std::runtime_error(std::string("can't open video file: " + m_file_name));
 662         }
 663         else if (m_camera_id != -1)
 664         {
 665             m_cap.open(m_camera_id);
 666             if (!m_cap.isOpened())
 667             {
 668                 std::stringstream msg;
 669                 msg << "can't open camera: " << m_camera_id;
 670                 throw std::runtime_error(msg.str());
 671             }
 672         }
 673         else
 674             throw std::runtime_error(std::string("specify video source"));
 675     }
 676 
 677     catch (std::exception e)
 678     {
 679         cerr << "ERROR: " << e.what() << std::endl;
 680         return -1;
 681     }
 682 
 683     return 0;
 684 } // initVideoSource()
 685 
 686 
 687 // this function is an example of "typical" OpenCL processing pipeline
 688 // It creates OpenCL buffer or image, depending on use_buffer flag,
 689 // from input media frame and process these data
 690 // (inverts each pixel value in half of frame) with OpenCL kernel
 691 int App::process_frame_with_open_cl(cv::Mat& frame, bool use_buffer, cl_mem* mem_obj)
 692 {
 693     cl_int res = CL_SUCCESS;
 694 
 695     CV_Assert(mem_obj);
 696 
 697     cl_kernel kernel = 0;
 698     cl_mem mem = mem_obj[0];
 699 
 700     if (0 == mem || 0 == m_img_src)
 701     {
 702         // allocate/delete cl memory objects every frame for the simplicity.
 703         // in real applicaton more efficient pipeline can be built.
 704 
 705         if (use_buffer)
 706         {
 707             cl_mem_flags flags = CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR;
 708 
 709             mem = clCreateBuffer(m_context, flags, frame.total(), frame.ptr(), &res);
 710             if (0 == mem || CL_SUCCESS != res)
 711                 return -1;
 712 
 713             res = clSetKernelArg(m_kernelBuf, 0, sizeof(cl_mem), &mem);
 714             if (CL_SUCCESS != res)
 715                 return -1;
 716 
 717             res = clSetKernelArg(m_kernelBuf, 1, sizeof(int), &frame.step[0]);
 718             if (CL_SUCCESS != res)
 719                 return -1;
 720 
 721             res = clSetKernelArg(m_kernelBuf, 2, sizeof(int), &frame.rows);
 722             if (CL_SUCCESS != res)
 723                 return -1;
 724 
 725             int cols2 = frame.cols / 2;
 726             res = clSetKernelArg(m_kernelBuf, 3, sizeof(int), &cols2);
 727             if (CL_SUCCESS != res)
 728                 return -1;
 729 
 730             kernel = m_kernelBuf;
 731         }
 732         else
 733         {
 734             cl_mem_flags flags_src = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR;
 735 
 736             cl_image_format fmt;
 737             fmt.image_channel_order     = CL_R;
 738             fmt.image_channel_data_type = CL_UNSIGNED_INT8;
 739 
 740             cl_image_desc desc_src;
 741             desc_src.image_type        = CL_MEM_OBJECT_IMAGE2D;
 742             desc_src.image_width       = frame.cols;
 743             desc_src.image_height      = frame.rows;
 744             desc_src.image_depth       = 0;
 745             desc_src.image_array_size  = 0;
 746             desc_src.image_row_pitch   = frame.step[0];
 747             desc_src.image_slice_pitch = 0;
 748             desc_src.num_mip_levels    = 0;
 749             desc_src.num_samples       = 0;
 750             desc_src.buffer            = 0;
 751             m_img_src = clCreateImage(m_context, flags_src, &fmt, &desc_src, frame.ptr(), &res);
 752             if (0 == m_img_src || CL_SUCCESS != res)
 753                 return -1;
 754 
 755             cl_mem_flags flags_dst = CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR;
 756 
 757             cl_image_desc desc_dst;
 758             desc_dst.image_type        = CL_MEM_OBJECT_IMAGE2D;
 759             desc_dst.image_width       = frame.cols;
 760             desc_dst.image_height      = frame.rows;
 761             desc_dst.image_depth       = 0;
 762             desc_dst.image_array_size  = 0;
 763             desc_dst.image_row_pitch   = 0;
 764             desc_dst.image_slice_pitch = 0;
 765             desc_dst.num_mip_levels    = 0;
 766             desc_dst.num_samples       = 0;
 767             desc_dst.buffer            = 0;
 768             mem = clCreateImage(m_context, flags_dst, &fmt, &desc_dst, 0, &res);
 769             if (0 == mem || CL_SUCCESS != res)
 770                 return -1;
 771 
 772             size_t origin[] = { 0, 0, 0 };
 773             size_t region[] = { (size_t)frame.cols, (size_t)frame.rows, 1 };
 774             res = clEnqueueCopyImage(m_queue, m_img_src, mem, origin, origin, region, 0, 0, &m_event);
 775             if (CL_SUCCESS != res)
 776                 return -1;
 777 
 778             res = clWaitForEvents(1, &m_event);
 779             if (CL_SUCCESS != res)
 780                 return -1;
 781 
 782             res = clSetKernelArg(m_kernelImg, 0, sizeof(cl_mem), &m_img_src);
 783             if (CL_SUCCESS != res)
 784                 return -1;
 785 
 786             res = clSetKernelArg(m_kernelImg, 1, sizeof(cl_mem), &mem);
 787             if (CL_SUCCESS != res)
 788                 return -1;
 789 
 790             kernel = m_kernelImg;
 791         }
 792     }
 793 
 794     m_event = clCreateUserEvent(m_context, &res);
 795     if (0 == m_event || CL_SUCCESS != res)
 796         return -1;
 797 
 798     // process left half of frame in OpenCL
 799     size_t size[] = { (size_t)frame.cols / 2, (size_t)frame.rows };
 800     res = clEnqueueNDRangeKernel(m_queue, kernel, 2, 0, size, 0, 0, 0, &m_event);
 801     if (CL_SUCCESS != res)
 802         return -1;
 803 
 804     res = clWaitForEvents(1, &m_event);
 805     if (CL_SUCCESS != res)
 806         return - 1;
 807 
 808     mem_obj[0] = mem;
 809 
 810     return  0;
 811 }
 812 
 813 
 814 // this function is an example of interoperability between OpenCL buffer
 815 // and OpenCV UMat objects. It converts (without copying data) OpenCL buffer
 816 // to OpenCV UMat and then do blur on these data
 817 int App::process_cl_buffer_with_opencv(cl_mem buffer, size_t step, int rows, int cols, int type, cv::UMat& u)
 818 {
 819     cv::ocl::convertFromBuffer(buffer, step, rows, cols, type, u);
 820 
 821     // process right half of frame in OpenCV
 822     cv::Point pt(u.cols / 2, 0);
 823     cv::Size  sz(u.cols / 2, u.rows);
 824     cv::Rect roi(pt, sz);
 825     cv::UMat uroi(u, roi);
 826     cv::blur(uroi, uroi, cv::Size(7, 7), cv::Point(-3, -3));
 827 
 828     if (buffer)
 829         clReleaseMemObject(buffer);
 830     m_mem_obj = 0;
 831 
 832     return 0;
 833 }
 834 
 835 
 836 // this function is an example of interoperability between OpenCL image
 837 // and OpenCV UMat objects. It converts OpenCL image
 838 // to OpenCV UMat and then do blur on these data
 839 int App::process_cl_image_with_opencv(cl_mem image, cv::UMat& u)
 840 {
 841     cv::ocl::convertFromImage(image, u);
 842 
 843     // process right half of frame in OpenCV
 844     cv::Point pt(u.cols / 2, 0);
 845     cv::Size  sz(u.cols / 2, u.rows);
 846     cv::Rect roi(pt, sz);
 847     cv::UMat uroi(u, roi);
 848     cv::blur(uroi, uroi, cv::Size(7, 7), cv::Point(-3, -3));
 849 
 850     if (image)
 851         clReleaseMemObject(image);
 852     m_mem_obj = 0;
 853 
 854     if (m_img_src)
 855         clReleaseMemObject(m_img_src);
 856     m_img_src = 0;
 857 
 858     return 0;
 859 }
 860 
 861 
 862 int App::run()
 863 {
 864     if (0 != initOpenCL())
 865         return -1;
 866 
 867     if (0 != initVideoSource())
 868         return -1;
 869 
 870     Mat img_to_show;
 871 
 872     // set running state until ESC pressed
 873     setRunning(true);
 874     // set process flag to show some data processing
 875     // can be toggled on/off by 'p' button
 876     setDoProcess(true);
 877     // set use buffer flag,
 878     // when it is set to true, will demo interop opencl buffer and cv::Umat,
 879     // otherwise demo interop opencl image and cv::UMat
 880     // can be switched on/of by SPACE button
 881     setUseBuffer(true);
 882 
 883     // Iterate over all frames
 884     while (isRunning() && nextFrame(m_frame))
 885     {
 886         cv::cvtColor(m_frame, m_frameGray, COLOR_BGR2GRAY);
 887 
 888         UMat uframe;
 889 
 890         // work
 891         timerStart();
 892 
 893         if (doProcess())
 894         {
 895             process_frame_with_open_cl(m_frameGray, useBuffer(), &m_mem_obj);
 896 
 897             if (useBuffer())
 898                 process_cl_buffer_with_opencv(
 899                     m_mem_obj, m_frameGray.step[0], m_frameGray.rows, m_frameGray.cols, m_frameGray.type(), uframe);
 900             else
 901                 process_cl_image_with_opencv(m_mem_obj, uframe);
 902         }
 903         else
 904         {
 905             m_frameGray.copyTo(uframe);
 906         }
 907 
 908         timerEnd();
 909 
 910         uframe.copyTo(img_to_show);
 911 
 912         putText(img_to_show, "Version : " + m_platformInfo.Version(), Point(5, 30), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
 913         putText(img_to_show, "Name : " + m_platformInfo.Name(), Point(5, 60), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
 914         putText(img_to_show, "Device : " + m_deviceInfo.Name(), Point(5, 90), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
 915         cv::String memtype = useBuffer() ? "buffer" : "image";
 916         putText(img_to_show, "interop with OpenCL " + memtype, Point(5, 120), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
 917         putText(img_to_show, "Time : " + timeStr() + " msec", Point(5, 150), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
 918 
 919         imshow("opencl_interop", img_to_show);
 920 
 921         handleKey((char)waitKey(3));
 922     }
 923 
 924     return 0;
 925 }
 926 
 927 
 928 void App::handleKey(char key)
 929 {
 930     switch (key)
 931     {
 932     case 27:
 933         setRunning(false);
 934         break;
 935 
 936     case ' ':
 937         setUseBuffer(!useBuffer());
 938         break;
 939 
 940     case 'p':
 941     case 'P':
 942         setDoProcess( !doProcess() );
 943         break;
 944 
 945     default:
 946         break;
 947     }
 948 }
 949 
 950 
 951 inline void App::timerStart()
 952 {
 953     m_t0 = getTickCount();
 954 }
 955 
 956 
 957 inline void App::timerEnd()
 958 {
 959     m_t1 = getTickCount();
 960     int64 delta = m_t1 - m_t0;
 961     m_time = (delta / m_frequency) * 1000; // units msec
 962 }
 963 
 964 
 965 inline string App::timeStr() const
 966 {
 967     stringstream ss;
 968     ss << std::fixed << std::setprecision(1) << m_time;
 969     return ss.str();
 970 }
 971 
 972 
 973 int main(int argc, char** argv)
 974 {
 975     const char* keys =
 976         "{ help h ?    |          | print help message }"
 977         "{ camera c    | -1       | use camera as input }"
 978         "{ video  v    |          | use video as input }";
 979 
 980     CommandLineParser cmd(argc, argv, keys);
 981     if (cmd.has("help"))
 982     {
 983         cmd.printMessage();
 984         return EXIT_SUCCESS;
 985     }
 986 
 987     App app(cmd);
 988 
 989     try
 990     {
 991         app.run();
 992     }
 993 
 994     catch (const cv::Exception& e)
 995     {
 996         cout << "error: " << e.what() << endl;
 997         return 1;
 998     }
 999 
1000     catch (const std::exception& e)
1001     {
1002         cout << "error: " << e.what() << endl;
1003         return 1;
1004     }
1005 
1006     catch (...)
1007     {
1008         cout << "unknown exception" << endl;
1009         return 1;
1010     }
1011 
1012     return EXIT_SUCCESS;
1013 } // main()
View Code

     官方给出的代码,有很多和OpenCL平台相关的处理部分,看起的比较冗杂,因此我们就不直接移植这个代码了(参考部分思想和API接口的基本调用),重新写一个新的基于Opencv3.1.0的OCL程序的Demo框架,基于TI的AM57x系列的板卡,其他的设备类似:

1. Opencv3.1版本OpenCL支持情况:

        首先对于Opencv3.1.0版本,从正式发布Opencv3开始,其对OpenCL的支持已经发生了很大的变化,在之前需要使用cv::ocl相关函数API来实现kernel的编译调用等等,在其中还包括了很多的数据搬移,而Opencv3正式改变了这样的情况,重新封装了一个新的数据类型cv::UMat,这个数据类型能够无缝对接Opencv的普通接口,从而最少的改动代码而最大的完成OpenCL平台的加速功能[OpenCV3.x-OpenCL.pptx]

         如上图所示,只需要将原来的Mat格式换为UMat格式就可以实现Opencv函数在OpenCL设备上加速运行,而这其中具体实施的基本原理是什么呢?接下来看一下其底层实现的基本原理,具体参看Opencv中OpenCL部分实现的源代码:

Reference:

官方代码说明:https://www.khronos.org/opencl/

Opencv(编译OCL模块)+VSC+MinGW环境搭建:https://www.cnblogs.com/uestc-mm/p/12758110.html

TI官方OpenCV3.1-Release:https://git.ti.com/cgit/opencv/tiopencv/?h=tiopencvrelease_3.1

TI官方AM57x Processor SDK Linux(OpenCV支持):http://software-dl.ti.com/processor-sdk-linux/esd/docs/latest/linux/Foundational_Components_OpenCV.html

标签:CL,cl,OpenCL,param,Opencv,DEVICE,query,id
来源: https://www.cnblogs.com/uestc-mm/p/12963792.html