使用OpenCL+OpenCV實現影象旋轉（一）
[題外話]近期申請了一個微信公眾號：平凡程式人生。有興趣的朋友可以關注，那裡將會涉及更多更新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> 呼叫OpenCVAPI imread()讀取一張彩色JPEG圖片，將它儲存在MAT變數中。該變數的data成員中儲存著將JPEG圖片解碼後的RGB資料。
2> 呼叫OpenCVAPI 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> 呼叫OpenCLAPI clGetPlatformIDs()直接獲取第一個可用的平臺資訊。該函式一般是先用它獲取支援OpenCL平臺的數目，然後再次呼叫它獲取某個平臺的資訊。兩次呼叫，通過傳遞不同引數區分。
6> 呼叫OpenCLAPI 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 APIclCreateProgramWithSource()，使用kernel的原始碼建立program物件。
14> 呼叫OpenCL APIclBuildProgram()編譯program物件。
15> 呼叫OpenCL APIclCreateKernel()，使用編譯完的程式物件建立kernel。
16> 呼叫OpenCL APIclSetKernelArg()為kernel程式傳遞引數，包括輸入輸出buffer地址，影象解析度和sin()\cos()值。
17> 呼叫OpenCL APIclEnqueueNDRangeKernel()執行kernel。
18> 呼叫OpenCL APIclEnqueueReadBuffer，將處理完的影象資料已經從裝置端傳遞到了host端的輸出buffer中。
19> 將輸出buffer中的資料copy到MAT物件的成員data中。
20> 呼叫OpenCV APIimwrite()將旋轉後的灰度影象儲存到檔案中，編碼為JPEG儲存起來。
21> 釋放輸入輸出buffer空間，釋放OpenCL建立的各個物件。
3、kernel程式程式碼
我們先看一下kernel程式。Kernel程式是每個workitem需要執行的，它需要儲存在以cl為字尾的檔案中，比如：ImageRotate.cl。
Kernel程式定義如下：
__kernel voidimg_rotate(
__global unsigned char*dest_data,
__global unsigned char*src_data,
int W,
int H,
floatsinTheta,
floatcosTheta)
有幾點需要注意的地方：
1〉 必須帶著關鍵字__kernel；
2〉 返回值必須為void；
3〉 區分清楚所傳引數的儲存型別，比如帶__global表示儲存在globalmemory中；什麼都不帶的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. }
(未完待續)