[题外话]近期申请了一个微信公众号：平凡程式人生。有兴趣的朋友可以关注，那里将会涉及更多更新OpenCL+OpenCV以及图像处理方面的文章。
最近在学习《OPENCL异构计算》，其中有一个实例是使用OpenCL实现图像旋转。这个实例中并没有涉及读取、保存、显示图像等操作，其中也存在一些小bug。在学习OpenCL之初，完整地实现这个实例还是很有意义的事情。
1、图像旋转原理
所谓图像旋转是指图像以某一点为中心旋转一定的角度，形成一幅新的图像的过程。这个点通常就是图像的中心。
由于是按照中心旋转，所以有这样一个属性：旋转前和旋转后的点离中心的位置不变.
根据这个属性，可以得到旋转后的点的坐标与原坐标的对应关系。
原图像的坐标一般是以左上角为原点的，我们先把坐标转换为以图像中心为原点。假设原图像的宽为w，高为h，（x0,y0）为原坐标内的一点，转换坐标后的点为（x1，y1）。可以得到：
X0’ = x0 - w/2;
y1’ = -y0 + h/2;
在新的坐标系下，假设点（x0，y0）距离原点的距离为r，点与原点之间的连线与x轴的夹角为b，旋转的角度为a，旋转后的点为（x1，y1）, 如下图所示。
那么有以下结论：
x0=r*cosb；y0=r*sinb
x1 = r*cos(b-a) = r*cosb*cosa+r*sinb*sina=x0*cosa+y0*sina；
y1=r*sin(b-a)=r*sinb*cosa-r*cosb*sina=-x0*sina+y0*cosa；
得到了转换后的坐标，我们只需要把这些坐标再转换为原坐标系即可。
x1’ = x1+w/2= x0*cosa+y0*sina+w/2
y1’=-y1+h/2=-(-x0*sina+y0*cosa)+h/2= x0*sina-y0*cosa+h/2
此处的x0/y0是新的坐标系中的值，转换为原坐标系为：
x1’ = x0*cosa+y0*sina+w/2=(x00-w/2)*consa+(-y00+h/2)*sina+w/2
y1’= x0*sina-y0*cosa+h/2=( x00-w/2)*sina-(-y00+h/2)*cosa+h/2
=(y00-h/2)*cosa+( x00-w/2)*sina+h/2
2、程序设计
对于图像旋转这个实例，为了处理简单，我将在灰度图上去做旋转。 大致的处理流程如下：
1> 调用OpenCV API imread()读取一张彩色JPEG图片，将它存储在MAT变量中。该变量的data成员中存储着将JPEG图片解码后的RGB数据。
2> 调用OpenCV API cvtColor()将存储RGB数据的MAT变量转换为只存储灰度图像数据的MAT对象。也可以使用函数imread()时直接将JPEG图像解码转换为灰度图像。
3> MAT对象的成员width和height存储着解码后图像的分辨率信息。根据当前分辨率，分配处理图像时所用的输入buffer和输出buffer。它们都按照存储char型数据进行空间申请。
4> 将MAT对象的成员data中数据copy到输入buffer中。同时将输出buffer初始化为全0。到此，我们调用OpenCV的API所要做的事情告一段落了。接下来就要调用OpenCL的API做事情了。
5> 调用OpenCL API clGetPlatformIDs()直接获取第一个可用的平台信息。该函数一般是先用它获取支持OpenCL平台的数目，然后再次调用它获取某个平台的信息。两次调用，通过传递不同参数区分。
6> 调用OpenCL API clGetDeviceIDs()获取第一个平台中第一个可用的设备。同样，这个函数也可以调用两次，分别获取当前平台的设备数目，再获取某个设备信息。
7> 调用OpenCL API clCreateContext()创建上下文。
8> 调用OpenCL API clCreateCommandQueue()创建host与device之间交互的command队列。
9> 调用OpenCL API clCreateBuffer()在设备端分配存储输入图像的buffer。
10> 调用OpenCL API clEnqueueWriteBuffer()将之前存储灰度图像数据的输入buffer内存copy到设备端buffer中。
11> 调用OpenCL API clCreateBuffer()在设备端分配处理完数据的存储buffer。
12> 调用文件读取函数，将kernel文件ImageRotate.cl中的内容读取到string变量中。
13> 调用OpenCL API clCreateProgramWithSource()，使用kernel的源码创建program对象。
14> 调用OpenCL API clBuildProgram()编译program对象。
15> 调用OpenCL API clCreateKernel()，使用编译完的程序对象创建kernel。
16> 调用OpenCL API clSetKernelArg()为kernel程序传递参数，包括输入输出buffer地址，图像分辨率和sin()\cos()值。
17> 调用OpenCL API clEnqueueNDRangeKernel()执行kernel。
18> 调用OpenCL API clEnqueueReadBuffer，将处理完的图像数据已经从设备端传递到了host端的输出buffer中。
19> 将输出buffer中的数据copy到MAT对象的成员data中。
20> 调用OpenCV API imwrite()将旋转后的灰度图像保存到文件中，编码为JPEG保存起来。
21> 释放输入输出buffer空间，释放OpenCL创建的各个对象。
3、kernel程序代码
我们先看一下kernel程序。Kernel程序是每个work item需要执行的，它需要存储在以cl为后缀的文件中，比如：ImageRotate.cl。
Kernel程序定义如下：
__kernel void img_rotate(
__global unsigned char *dest_data,
__global unsigned char *src_data,
int W,
int H,
float sinTheta,
float cosTheta)
有几点需要注意的地方：
1〉 必须带着关键字__kernel；
2〉 返回值必须为void；
3〉 区分清楚所传参数的存储类型，比如带__global表示存储在global memory中；什么都不带的W、H等表示存储在work item的private memory中。
Kernel程序如下：
1. __kernel void img_rotate(
2. __global unsigned char *dest_data,
3. __global unsigned char *src_data,
4. int W,
5. int H,
6. float sinTheta,
7. float cosTheta){
8. //work item gets its index within index space
9. const int ix = get_global_id(0);
10. const int iy = get_global_id(1);
11.
12. //calculate location of data to move int (ix, iy)
13. //output decomposition as mentioned
14. float xpos = ((float)(ix - W / 2)) * cosTheta + ((float)(-iy + H / 2)) * sinTheta + W / 2;
15. float ypos = ((float)(ix - W / 2)) * sinTheta + ((float)(iy - H / 2)) * cosTheta + H / 2;
16.
17. //bound checking
18. if (((int)xpos >=0) && ((int)xpos < W) &&
19. ((int)ypos >= 0) && ((int)ypos < H)) {
20. dest_data[(int)ypos * W + (int)xpos] = src_data[iy * W + ix];
21. }
22. }
(未完待续)
