OpenCL + OpenCV 图像旋转

▶ 使用 OpenCV 从文件读取彩色的 png 图像，旋转一定角度以后写回文件

● 代码，核函数

 1 // rotate.cl                                  
 2 //__constant  sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP;// 设备采样器，可以启用，并删除函数 imageRotate 中的采样器参数
 3 
 4 __kernel void imageRotate(__read_only image2d_t inputImage, __write_only image2d_t outputImage, float angle, sampler_t sampler)
 5 {
 6     const int width = get_image_width(inputImage), height = get_image_height(inputImage);
 7     const int halfWidth = width / 2, halfHeight = height / 2;
 8     const int x = get_global_id(0), y = get_global_id(1);
 9     const int xt = x - halfWidth, yt = y - halfHeight;
10     const float sinFactor = sin(angle), cosFactor = cos(angle);
11 
12     float2 readCoord = (float2)(halfWidth + cosFactor * xt - sinFactor * yt, readCoord.y = halfHeight + sinFactor * xt + cosFactor * yt);
13     float4 value = read_imagef(inputImage, sampler, readCoord);
14     write_imagef(outputImage, (int2)(x, y), value);
15     return;
16 }

● 代码，分三通道分别旋转

  1 #include <stdio.h>
  2 #include <stdlib.h>
  3 #include <cl.h>
  4 #include <opencv.hpp>
  5 #include <opencv2corecvstd.hpp>   // namespace cv 的定义
  6 
  7 #pragma warning(disable : 4996)     // 解封OPenCL1.2
  8 
  9 using namespace cv;
 10 
 11 const char *sourceProgram = "D:/Code/OpenCL/rotate.cl";// 核函数文件
 12 const char *imagePath = "D:\input.png";
 13 const float angle = 3.14f / 4;
 14 
 15 int readSource(const char* kernelPath, char **output)// 读取文本文件，存储为 char *
 16 {
 17     FILE *fp;
 18     int size;
 19     fopen_s(&fp, kernelPath, "rb");
 20     if (!fp)
 21     {
 22         printf("Open kernel file failed
");
 23         exit(-1);
 24     }
 25     if (fseek(fp, 0, SEEK_END) != 0)
 26     {
 27         printf("Seek end of file faildd
");
 28         exit(-1);
 29     }
 30     if ((size = ftell(fp)) < 0)
 31     {
 32         printf("Get file position failed
");
 33         exit(-1);
 34     }
 35     rewind(fp);
 36     if ((*output = (char *)malloc(size + 1)) == NULL)
 37     {
 38         printf("Allocate space failed
");
 39         exit(-1);
 40     }    
 41     fread((void*)*output, 1, size, fp);
 42     fclose(fp);
 43     (*output)[size] = '';
 44     printf("readSource succeed, program file: %s
", kernelPath);
 45     return size;
 46 }
 47 
 48 int main()
 49 {
 50     // 准备平台，设备，上下文，命令队列部分    
 51     cl_int status;
 52     cl_uint nPlatform;
 53     clGetPlatformIDs(0, NULL, &nPlatform);
 54     cl_platform_id *listPlatform = (cl_platform_id*)malloc(nPlatform * sizeof(cl_platform_id));
 55     clGetPlatformIDs(nPlatform, listPlatform, NULL);
 56     cl_uint nDevice = 0;
 57     clGetDeviceIDs(listPlatform[0], CL_DEVICE_TYPE_ALL, 0, NULL, &nDevice);
 58     cl_device_id *listDevice = (cl_device_id*)malloc(nDevice * sizeof(cl_device_id));
 59     clGetDeviceIDs(listPlatform[0], CL_DEVICE_TYPE_ALL, nDevice, listDevice, NULL);
 60     cl_context context = clCreateContext(NULL, nDevice, listDevice, NULL, NULL, &status);    
 61     cl_command_queue queue = clCreateCommandQueue(context, listDevice[0], 0, &status);                          // OpenCL1.2
 62     //cl_command_queue_properties queueProp[5] = {CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT,// OpenCL2.0
 63     //                                            CL_QUEUE_SIZE, CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE,
 64     //                                            0};
 65     //cl_command_queue queue = clCreateCommandQueueWithProperties(context, listDevice[0], &queueProp, &status); // 第三个参数 queueProp 各种改都会报内存越界 0xC0000005
 66 
 67     // 图片相关
 68     Mat image, channel[3];
 69     image = imread(imagePath);
 70     split(image, channel);                          // 拆分为三通道，分别旋转后拼合
 71     const int imageHeight = image.rows, imageWidth = image.cols;
 72     unsigned char *imageData = (unsigned char*)malloc(sizeof(unsigned char) * imageHeight * imageWidth);
 73     
 74     cl_image_format format;
 75     format.image_channel_order = CL_R;              // 单通道
 76     format.image_channel_data_type = CL_UNORM_INT8; // 无符号 8 位整形，0 ~ 255
 77     cl_image_desc desc;
 78     desc.image_type = CL_MEM_OBJECT_IMAGE2D;        // 可以 memset(desc,sizeof(cl_image_desc)); 后仅对前三项赋值
 79     desc.image_width = imageWidth;
 80     desc.image_height = imageHeight;
 81     desc.image_depth = 0;
 82     desc.image_array_size = 0;
 83     desc.image_row_pitch = 0;
 84     desc.image_slice_pitch = 0;
 85     desc.num_mip_levels = 0;
 86     desc.num_samples = 0;
 87     desc.buffer = NULL;
 88     cl_mem d_inputImage = clCreateImage(context, CL_MEM_READ_ONLY, &format, &desc, NULL, &status);
 89     cl_mem d_outputImage = clCreateImage(context, CL_MEM_WRITE_ONLY, &format, &desc, NULL, &status);
 90 
 91     // 采样器
 92     cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &status);  // OpenCL1.2
 93     //cl_sampler_properties samplerProp[7] = {CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE,                               // OpenCL2.0
 94     //                                        CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, 
 95     //                                        CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST,
 96     //                                        0};             
 97     //cl_sampler sampler = clCreateSamplerWithProperties(context, samplerProp, &status);                            // 也是内存越界，用不了
 98 
 99     // 程序和内核
100     char* source = NULL;
101     const size_t lenSource = readSource(sourceProgram, &source);
102     cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source, &lenSource, &status);
103     clBuildProgram(program, 1, listDevice, NULL, NULL, NULL);
104     cl_kernel kernel = clCreateKernel(program, "imageRotate", &status);    
105     clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_inputImage);
106     clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_outputImage);
107     clSetKernelArg(kernel, 2, sizeof(cl_float), &angle);
108     clSetKernelArg(kernel, 3, sizeof(cl_sampler), &sampler);
109     size_t origin[3] = { 0, 0, 0 }, region[3] = { imageWidth, imageHeight, 1 };// 拷贝图片缓冲区时使用的起点和范围参数
110     size_t globalSize[2] = { imageWidth, imageHeight };
111 
112     for (int i = 0; i < 3; i++)// 分三个通道拷入缓冲区，执行旋转操作，拷回内存
113     {
114         memcpy(imageData, channel[i].data, sizeof(unsigned char) * imageHeight * imageWidth);
115         clEnqueueWriteImage(queue, d_inputImage, CL_TRUE, origin, region, 0, 0, imageData, 0, NULL, NULL);
116         clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0, NULL, NULL);
117         clEnqueueReadImage(queue, d_outputImage, CL_TRUE, origin, region, 0, 0, imageData, 0, NULL, NULL);
118         memcpy(channel[i].data, imageData, sizeof(unsigned char) * imageHeight * imageWidth);
119     }
120     
121     merge(channel, 3, image);// 合并通道，结果写入文件，在窗口中展示结果
122     imwrite("D:/output.png", image);
123     imshow("Result", image);
124     waitKey(0);
125 
126     free(listPlatform);
127     free(listDevice);
128     clReleaseContext(context);
129     clReleaseMemObject(d_inputImage);
130     clReleaseMemObject(d_outputImage);
131     clReleaseCommandQueue(queue);
132     clReleaseProgram(program);
133     clReleaseKernel(kernel);    
134     //getchar();
135     return 0;
136 }

● 代码，四个通道同时操作，注意图片读入和输出的时候只有三个通道，需要进行调整

  1 #include <stdio.h>
  2 #include <stdlib.h>
  3 #include <cl.h>
  4 #include <opencv.hpp>
  5 #include <opencv2corecvstd.hpp>   // namespace cv 的定义
  6 
  7 #pragma warning(disable : 4996)     // 解封OPenCL1.2
  8 
  9 using namespace cv;
 10 
 11 const char *sourceProgram = "D:/Code/OpenCL/rotate.cl";// 核函数文件
 12 const char *imagePath = "D:/input.png";
 13 const float angle = 3.14f / 4;
 14 
 15 int readSource(const char* kernelPath, char **output)// 读取文本文件，存储为 char *
 16 {
 17     FILE *fp;
 18     int size;
 19     fopen_s(&fp, kernelPath, "rb");
 20     if (!fp)
 21     {
 22         printf("Open kernel file failed
");
 23         exit(-1);
 24     }
 25     if (fseek(fp, 0, SEEK_END) != 0)
 26     {
 27         printf("Seek end of file faildd
");
 28         exit(-1);
 29     }
 30     if ((size = ftell(fp)) < 0)
 31     {
 32         printf("Get file position failed
");
 33         exit(-1);
 34     }
 35     rewind(fp);
 36     if ((*output = (char *)malloc(size + 1)) == NULL)
 37     {
 38         printf("Allocate space failed
");
 39         exit(-1);
 40     }    
 41     fread((void*)*output, 1, size, fp);
 42     fclose(fp);
 43     (*output)[size] = '';
 44     printf("readSource succeed, program file: %s
", kernelPath);
 45     return size;
 46 }
 47 
 48 int main()
 49 {
 50     // 准备平台，设备，上下文，命令队列部分    
 51     cl_int status;
 52     cl_uint nPlatform;
 53     clGetPlatformIDs(0, NULL, &nPlatform);
 54     cl_platform_id *listPlatform = (cl_platform_id*)malloc(nPlatform * sizeof(cl_platform_id));
 55     clGetPlatformIDs(nPlatform, listPlatform, NULL);
 56     cl_uint nDevice = 0;
 57     clGetDeviceIDs(listPlatform[0], CL_DEVICE_TYPE_ALL, 0, NULL, &nDevice);
 58     cl_device_id *listDevice = (cl_device_id*)malloc(nDevice * sizeof(cl_device_id));
 59     clGetDeviceIDs(listPlatform[0], CL_DEVICE_TYPE_ALL, nDevice, listDevice, NULL);
 60     cl_context context = clCreateContext(NULL, nDevice, listDevice, NULL, NULL, &status);    
 61     cl_command_queue queue = clCreateCommandQueue(context, listDevice[0], 0, &status);// OpenCL1.2    
 62 
 63     // 图片相关
 64     Mat image = imread(imagePath);
 65     const int imageHeight = image.rows, imageWidth = image.cols;
 66     unsigned char *imageData = (unsigned char*)malloc(sizeof(unsigned char) * imageHeight * imageWidth * 4);
 67     
 68     for (int i = 0; i < imageWidth * imageHeight; i++)// imread 读进来只有 RGB 三个通道（可能跟图片本身有关），要补成 4 个通道
 69     {
 70         imageData[4 * i + 0] = image.data[3 * i + 2];//R
 71         imageData[4 * i + 1] = image.data[3 * i + 1];//G
 72         imageData[4 * i + 2] = image.data[3 * i + 0];//B
 73         imageData[4 * i + 3] = 255;                  //A
 74     }
 75 
 76     cl_image_format format;
 77     format.image_channel_order = CL_RGBA;            // 合并通道
 78     format.image_channel_data_type = CL_UNORM_INT8; // 无符号 8 位整形，0 ~ 255
 79     cl_image_desc desc;
 80     desc.image_type = CL_MEM_OBJECT_IMAGE2D;        // 可以 memset(desc,sizeof(cl_image_desc)); 后仅对前三项赋值
 81     desc.image_width = imageWidth;
 82     desc.image_height = imageHeight;
 83     desc.image_depth = 0;
 84     desc.image_array_size = 0;
 85     desc.image_row_pitch = 0;
 86     desc.image_slice_pitch = 0;
 87     desc.num_mip_levels = 0;
 88     desc.num_samples = 0;
 89     desc.buffer = NULL;
 90     cl_mem d_inputImage = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, &format, &desc, imageData, &status);// 输入图片直接在主机上
 91     cl_mem d_outputImage = clCreateImage(context, CL_MEM_WRITE_ONLY, &format, &desc, NULL, &status);    
 92 
 93     // 采样器
 94     cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &status);  // OpenCL1.2
 95 
 96     // 程序和内核
 97     char* source = NULL;
 98     const size_t lenSource = readSource(sourceProgram, &source);
 99     cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source, &lenSource, &status);
100     clBuildProgram(program, 1, listDevice, NULL, NULL, NULL);
101     cl_kernel kernel = clCreateKernel(program, "imageRotate", &status);    
102     clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_inputImage);
103     clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_outputImage);
104     clSetKernelArg(kernel, 2, sizeof(cl_float), &angle);
105     clSetKernelArg(kernel, 3, sizeof(cl_sampler), &sampler);
106     size_t origin[3] = { 0, 0, 0 }, region[3] = { imageWidth, imageHeight, 1 };// 拷贝图片缓冲区时使用的起点和范围参数
107     size_t globalSize[2] = { imageWidth, imageHeight };
108 
109     clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0, NULL, NULL);
110     clEnqueueReadImage(queue, d_outputImage, CL_TRUE, origin, region, 0, 0, imageData, 0, NULL, NULL);
111 
112     for (int i = 0; i < imageWidth * imageHeight; i++)// 去掉第 4 个通道，返回 image 中
113     {
114         image.data[3 * i + 0] = imageData[4 * i + 2];//B
115         image.data[3 * i + 1] = imageData[4 * i + 1];//G
116         image.data[3 * i + 2] = imageData[4 * i + 0];//R        
117     }    
118     
119     imwrite("D:/output.png", image);
120     imshow("Result", image);
121     waitKey(0);
122 
123     free(listPlatform);
124     free(listDevice);
125     clReleaseContext(context);
126     clReleaseMemObject(d_inputImage);
127     clReleaseMemObject(d_outputImage);
128     clReleaseCommandQueue(queue);
129     clReleaseProgram(program);
130     clReleaseKernel(kernel);    
131     //getchar();
132     return 0;
133 }

● 输入、输出结果，顺时针转 45 度，因为使用了最近邻采样，结果中锯齿比较严重

● 另一种解封旧 API 的方法，在包含头文件 <cl.h> 前使用 #define CL_USE_DEPRECATED_OPENCL_1_2_APIS ，其中 1_2 可以改成 1_0，1_1 等（https://stackoverflow.com/questions/28500496/opencl-function-found-deprecated-by-visual-studio/28500846#28500846）

● 使用 cl_command_queue_properties 和函数 clCreateCommandQueueWithProperties 来创建命令队列，或是用 cl_sampler_properties 和函数 clCreateSamplerWithProperties 来创建采样器都失败了，报内存访问越界错误（0xC0000005），无论是按格式书写还是把 queueProp 改成 0，创建时第三个参数写成 &queueProp 都不行；有人说更新显卡驱动以后就好了（https://stackoverflow.com/questions/39864947/opencl-cl-out-of-host-memory-on-clcreatecommandqueuewithproperties-with-minima）。最后解决了，用 AMD APP SDK 下面的动态库 amdocl64.dll 替换掉 C:WindowsSystem32 里边那个相同库就好了，可以完全使用 OpenCL2.0 的 API，不再报错。

● cv::imread 读入的图片是按照 [ R, G, B, R, G, B, R, G, B, ...] 存放的，在用 OpenCL处理之前需要进行一定的预处理，要么用 split 分解各通道为单独的图片，要么手工拆解，算完以后也要按照这种存放方式转回图像数据中。在发现通道个数和顺序的问题前，要么在调用函数 clCreateImage 的时候返回 -37，-38，-39，要么直接旋转得到像下面这样的图片。以后记得，如果出现这种交叉条纹的图像，有可能是通道交错导致的。

● 吐槽一下，网上能找到的 OpenCL + OpenCV 做图片旋转的基本上有几个版本（https://blog.csdn.net/c602273091/article/details/45418223，https://blog.csdn.net/icamera0/article/details/71598323，https://blog.csdn.net/jaccen2012/article/details/51367388）都是用 FreeImage 库把图像处理成灰度图来旋转的（参考了刘文志等(2016). OpenCL 异构并行计算[M]. 的代码？），输出肯定是灰度图了，然后大家博客就相互抄吧，全是垃圾。好不容易找到一个彩色的（https://blog.csdn.net/Bob_Dong/article/details/64906734）代码还看不了。