OPENCV 球面投影 opencv图像的旋转角度计算

转载

angel 2024-03-28 21:18:46

文章标签 OPENCV 球面投影算法 OpenCL CL 图像旋转 文章分类 计算机视觉人工智能

OpenCL笔记一

1. 图像旋转原理
2. OpenCL编程详细解析
3. 模块分析
4. Platform
5. 在 context 上查询 device
6. Running time
7. 加载 OpenCL 内核程序并创建一个 program 对象

先以图像旋转的实例，具体介绍OpenCL编程的步骤。首先给出实现流程，然后给出实现图像旋转的C循环实现和OpenCL C kernel实现。

1. 图像旋转原理

图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度，通常是指绕图像的中心以逆时针方向旋转。假设图像的左上角为(l, t), 右下角为(r, b)，则图像上任意点(x, y) 绕其中心(xcenter, ycenter)逆时针旋转θ角度后，新的坐标位置(x’,y’)的计算公式为：

x′ = (x - xcenter) cosθ - (y － ycenter) sinθ + xcenter,

y′ = (x - xcenter) sinθ + (y － ycenter) cosθ + ycenter.

C代码：

void rotate(
      unsigned char* inbuf,
      unsigned char* outbuf,
      int w, int h,
      float sinTheta,
      float cosTheta)
{
   int i, j;
   int xc = w/2;
   int yc = h/2;
   for(i = 0; i < h; i++)
   {
     for(j=0; j< w; j++)
     {
       int xpos =  (j-xc)*cosTheta - (i - yc) * sinTheta + xc;
       int ypos =  (j-xc)*sinTheta + (i - yc) * cosTheta + yc;
       if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)
          outbuf[ypos*w + xpos] = inbuf[i*w+j];
     }
   }
}

OpenCL C kernel代码：

#pragma OPENCL EXTENSION cl_amd_printf : enable
__kernel  void image_rotate(
      __global uchar * src_data,
      __global uchar * dest_data,        //Data in global memory
      int W,    int H,                   //Image Dimensions
      float sinTheta, float cosTheta )   //Rotation Parameters
{
   const int ix = get_global_id(0);
   const int iy = get_global_id(1);
   int xc = W/2;
   int yc = H/2;
   int xpos =  ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;
   int ypos =  (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;
   if ((xpos>=0) && (xpos< W)   && (ypos>=0) && (ypos< H))
      dest_data[ypos*W+xpos]= src_data[iy*W+ix];
}

正如上面代码中所给出的那样，在C代码中需要两重循环来计算横纵坐标上新的坐标位置。其实，在图像旋转的算法中每个点的计算可以独立进行，与其它点的坐标位置没有关系，所以并行处理较为方便。OpenCL C kernel代码中用了并行处理。

上面的代码在Intel的OpenCL平台上进行了测试，处理器为双核处理器，图像大小为4288*3216，如果用循环的方式运行时间稳定在0.256s左右，而如果用OpenCL C kernel并行的方式，运行时间稳定在0.132秒左右。GPU的测试在NVIDIA的GeForce G105M显卡上进行，运行时间稳定在0.0810s左右。从循环的方式，双核CPU并行以及GPU并行计算已经可以看出，OpenCL编程的确能大大提高执行效率。

2. OpenCL编程详细解析

OpenCL作为一门开源的异构并行计算语言，设计之初就是使用一种模型来模糊各种硬件差异。作为软件开发人员，关注的就是编程模型。OpenCL程序的流程大致如下：

· Platform

· 查询并选择一个 platform

· 在 platform 上创建 context

· 在 context 上查询并选择一个或多个 device

· Running time

· 加载 OpenCL 内核程序并创建一个 program 对象

· 为指定的 device 编译 program 中的 kernel

· 创建指定名字的 kernel 对象

· 为 kernel 创建内存对象

· 为 kernel 设置参数

· 在指定的 device 上创建 command queue

· 将要执行的 kernel 放入 command queue

· 将结果读回 host

· 资源回收

3. 模块分析

使用 OpenCL API 编程与一般 C/C++ 引入第三方库编程没什么区别。所以，首先要做的自然是 include 相关的头文件。由于在 MacOS X 10.6下OpenCL的头文件命名与其他系统不同，通常使用一个#if defined进行区分，代码如下：

#if defined(__APPLE__) || defined(__MACOSX)
  #include <OpenCL/cl.hpp>
#else
  #include <CL/cl.h>
#endif

接下来就进入真正的编码流程了。

4. Platform

查询并选择一个 platform
首先要取得系统中所有的 OpenCL platform。所谓的 platform 指的就是硬件厂商提供的 OpenCL 框架，不同的 CPU/GPU 开发商（比如 Intel、AMD、Nvdia）可以在一个系统上分别定义自己的 OpenCL 框架。所以需要查询系统中可用的 OpenCL 框架，即 platform。使用 API 函数 clGetPlatformIDs 获取可用 platform 的数量：

cl_int status = 0;
  cl_uint numPlatforms;
  cl_platform_id platform = NULL;
  status = clGetPlatformIDs( 0, NULL, &numPlatforms);
   
  if(status != CL_SUCCESS){
      printf("Error: Getting Platforms\n");
      return EXIT_FAILURE;
  }

然后根据数量来分配内存，并得到所有可用的 platform，所使用的 API 还是clGetPlatformIDs。在 OpenCL 中，类似这样的函数调用很常见：第一次调用以取得数目，便于分配足够的内存；然后调用第二次以获取真正的信息。

if (numPlatforms > 0) {
      cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
      status = clGetPlatformIDs(numPlatforms, platforms, NULL);
     if (status != CL_SUCCESS) {
         printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
          return -1;
     }

现在，所有的 platform 都存在了变量 platforms 中，接下来需要做的就是取得所需的 platform。本人的PC上配置的是 Intel 处理器和 AMD 显卡，专业点的说法叫 Intel 的 CPU 和 NVIDIA的 GPU 😃。所以这儿有两套 platform，为了体验下 GPU 的快感，所以使用 AMD 的 platform。通过使用 clGetPlatformInfo 来获得 platform 的信息。通过这个 API 可以知晓 platform 的厂商信息，以便选出需要的 platform。代码如下：

for (unsigned int i = 0; i < numPlatforms; ++i) {
          char pbuff[100];
          status = clGetPlatformInfo(
                       platforms[i],
                       CL_PLATFORM_VENDOR,
                      sizeof(pbuff),
                      pbuff,
                      NULL);
          platform = platforms[i];
         if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
             break;
         }
     }

不同的厂商信息可以参考 OpenCL Specifications，这儿只是简单的筛选出 AMD 。

在 platform 上建立 context
第一步是通过 platform 得到相应的 context properties

// 如果能找到相应平台，就使用，否则返回NULL
  cl_context_properties cps[3] = {
      CL_CONTEXT_PLATFORM,
      (cl_context_properties)platform,
      0
  };
   
  cl_context_properties *cprops = (NULL == platform) ? NULL : cps;

第二步是通过 clCreateContextFromType 函数创建 context。

// 生成 context
  cl_context context = clCreateContextFromType(
                           cprops,
                          CL_DEVICE_TYPE_GPU,
                          NULL,
                          NULL,
                          &status);
 if (status != CL_SUCCESS) {
      printf("Error: Creating Context.(clCreateContexFromType)\n");
    return EXIT_FAILURE;
 }

函数的第二个参数可以设定 context 关联的设备类型。本例使用的是 GPU 作为OpenCL计算设备。目前可以使用的类别包括：

- CL_DEVICE_TYPE_CPU
- CL_DEVICE_TYPE_GPU
- CL_DEVICE_TYPE_ACCELERATOR
- CL_DEVICE_TYPE_DEFAULT
- CL_DEVICE_TYPE_ALL

5. 在 context 上查询 device

context 创建好之后，要做的就是查询可用的 device。

status = clGetContextInfo(context,
                            CL_CONTEXT_DEVICES,
                           0,
                           NULL,
                            &deviceListSize);
  if (status != CL_SUCCESS) {
      printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
      return EXIT_FAILURE;
  }
 cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
 if (devices == 0) {
     printf("Error: No devices found.\n");
    return EXIT_FAILURE;
 }
  
 status = clGetContextInfo(context,
                          CL_CONTEXT_DEVICES,
                           deviceListSize,
                           devices,
                           NULL);
 if (status != CL_SUCCESS) {
     printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
     return EXIT_FAILURE;
 }

与获取 platform 类似，调用两次 clGetContextInfo 来完成查询。第一次调用获取关联 context 的 device 个数，并根据个数申请内存；第二次调用获取所有 device 实例。如果想了解每个 device 的具体信息，可以调用 clGetDeviceInfo 函数来获取，返回的信息有设备类型、生产商以及设备对某些扩展功能的支持与否等等。详细使用情况请参阅 OpenCL Specifications。

到此，platform 相关的程序已经准备就绪了，下面到此的完整代码：

1.  /* OpenCL_01.cpp 
2.   * (c) by keyring <keyrings@163.com>
3.   * 2013.10.26
4.   */
5.   
6.  #if defined(__APPLE__) || defined(__MACOSX)
7.  #include <OpenCL/cl.hpp>
8.  #else
9.  #include <CL/cl.h>
10. #endif
11.  
12. #include <iostream>
13.  
14. int main(int argc, char const *argv[])
15. {
16.     printf("hello OpenCL\n");
17.     cl_int status = 0;
18.     size_t deviceListSize;
19.  
20.     // 得到并选择可用平台
21.     cl_uint numPlatforms;
22.     cl_platform_id platform = NULL;
23.     status = clGetPlatformIDs(0, NULL, &numPlatforms);
24.  
25.     if (status != CL_SUCCESS) {
26.         printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");
27.         return EXIT_FAILURE;
28.     }
29.  
30.     if (numPlatforms > 0) {
31.         cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));
32.         status = clGetPlatformIDs(numPlatforms, platforms, NULL);
33.         if (status != CL_SUCCESS) {
34.             printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");
35.             return -1;
36.         }
37.  
38.         // 遍历所有 platform，选择想用的
39.         for (unsigned int i = 0; i < numPlatforms; ++i) {
40.             char pbuff[100];
41.             status = clGetPlatformInfo(
42.                          platforms[i],
43.                          CL_PLATFORM_VENDOR,
44.                          sizeof(pbuff),
45.                          pbuff,
46.                          NULL);
47.             platform = platforms[i];
48.             if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {
49.                 break;
50.             }
51.         }
52.  
53.         delete platforms;
54.     }
55.  
56.     // 如果能找到相应平台，就使用，否则返回NULL
57.     cl_context_properties cps[3] = {
58.         CL_CONTEXT_PLATFORM,
59.         (cl_context_properties)platform,
60.         0
61.     };
62.  
63.     cl_context_properties *cprops = (NULL == platform) ? NULL : cps;
64.  
65.  
66.     // 生成 context
67.     cl_context context = clCreateContextFromType(
68.                              cprops,
69.                              CL_DEVICE_TYPE_GPU,
70.                              NULL,
71.                              NULL,
72.                              &status);
73.     if (status != CL_SUCCESS) {
74.         printf("Error: Creating Context.(clCreateContexFromType)\n");
75.         return EXIT_FAILURE;
76.     }
77.  
78.     // 寻找OpenCL设备
79.  
80.     // 首先得到设备列表的长度
81.     status = clGetContextInfo(context,
82.                               CL_CONTEXT_DEVICES,
83.                               0,
84.                               NULL,
85.                               &deviceListSize);
86.     if (status != CL_SUCCESS) {
87.         printf("Error: Getting Context Info device list size, clGetContextInfo)\n");
88.         return EXIT_FAILURE;
89.     }
90.     cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);
91.     if (devices == 0) {
92.         printf("Error: No devices found.\n");
93.         return EXIT_FAILURE;
94.     }
95.  
96.     // 然后得到设备列表
97.     status = clGetContextInfo(context,
98.                               CL_CONTEXT_DEVICES,
99.                               deviceListSize,
100.                                        devices,
101.                                        NULL);
102.              if (status != CL_SUCCESS) {
103.                  printf("Error: Getting Context Info (device list, clGetContextInfo)\n");
104.                  return EXIT_FAILURE;
105.              }

6. Running time

前面写了这么多，其实还没真正进入具体的程序逻辑中，顶多算配好了 OpenCL 运行环境。真正的逻辑代码，即程序的任务就是运行时模块。本例的任务是在一个 4×4的二维空间上，按一定的规则给每个元素赋值，具体代码如下：

#define KERNEL(...)#__VA_ARGS__
   
  const char *kernelSourceCode = KERNEL(
                                     __kernel void hellocl(__global uint *buffer)
  {
      size_t gidx = get_global_id(0);
      size_t gidy = get_global_id(1);
      size_t lidx = get_local_id(0);
      buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);
  
 }
                               );

这一段就是真正的逻辑，也就是代码要干的事。使用的是 OpenCL 自定的一门类C语言，具体的语法什么的现在先不纠结。这段代码是直接嵌入 cpp 文件的静态字符串。也可以将 kernel 程序单独写成一个文件。