OpenCL 第7課:旋轉變換(1)

旋轉是一個常用的處理功能。圖片中所有的點以某一個點為軸，順時或逆時方向旋轉N個角度。我們利用OpenCL就可以對圖片中所有的點進行并行轉換，大大提高效率。

上兩節中，我們編寫了CL文件來傳遞數組的地址，這一節中我們會多加入幾個參數傳遞。

首先我們先來看下圖片旋轉的原理。這里我們假設圖片的旋轉是以圖片的中心點為軸。也就是（width/2,height/2）這個點。旋轉的角度是任意值。圖片旋轉會出現這兩種情況。

TM截圖20130327230153

第一幅是原圖，第二幅是旋轉了30度，大家可以看到旋轉后圖片的一部份數據已經超出了原來圖片的大小范圍。怎么處理超出部份的數據呢。因為我們的圖片數據是用數組來存儲的，一般我們有兩個方法，1、存儲圖片的數組做大些。2、超出圖片部份的數據不顯示。這里我們選擇后者，方便。

下面我們來看下rotate.cl程序和主程序。

rotate.cl程序

?
__kernel void rotation(__global int* A,
????????????????????__global int* B,
????????????????????int width,
????????????????????int height,
????????????????????float sinangle,
????????????????????float cosangle)
{
????//獲取索引號，這里是二維的，所以可以取兩個
????//否則另一個永遠是0
????int col = get_global_id(0);
????int row = get_global_id(1);
????//計算圖形中心點
????float cx = ((float)width)/2;
????float cy = ((float)height)/2;
????int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
????int ny = (int)(cy + (-1*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy));
????//邊界檢測
????if(nx>=0 && nx<width && ny>=0 && ny<height)
????{
????????B[nx + ny*width] = A[col + row*width];
????}
}

主程序

?
#include <iostream>
#include <stdio.h>
#include <string.h>
#include <string>
#include <conio.h>
#include <math.h>//數學庫
#include <CL/cl.h>//包含CL的頭文件
using namespace std;
//8x8數組
const int dim_x = 8;
const int dim_y = 8;
//45度的弧度
const float angle = 3.1415926f/4.0f;
//從外部文件獲取cl內核代碼
bool GetFileData(const char* fname,string& str)
{
????FILE* fp = fopen(fname,"r");
????if(fp==NULL)
????{
????????printf("no found file\n");
????????return false;
????}
????while(feof(fp)==0)
????{
????????str += fgetc(fp);
????}
????return true;
}
int  main()
{
????//先讀外部CL核心代碼，如果失敗則退出。
????//代碼存buf_code里面
????string code_file;
????if(false == GetFileData("rotate.cl",code_file))
????{
????????printf("Open rotate.cl error\n");
????????return 0;
????}
????char* buf_code = new char[code_file.size()];
????strcpy(buf_code,code_file.c_str());
????buf_code[code_file.size()-1] = NULL;
????//聲明CL所需變量。
????cl_device_id device;
????cl_platform_id platform_id = NULL;
????cl_context context;
????cl_command_queue cmdQueue;
????cl_mem bufferA,bufferB;
????cl_program program;
????cl_kernel kernel = NULL;
????//我們使用的是二維向量
????//設定向量大小（維數）
????size_t globalWorkSize[2];
????globalWorkSize[0] = dim_x ;
????globalWorkSize[1] = dim_y;
????cl_int err;
????/*
????????定義輸入變量和輸出變量，并設定初值
????*/
????size_t datasize = sizeof(int) * dim_x * dim_y;
????int m,n;
????int buf_A[] = {0,0,0,1,1,0,0,0,
???????????????????0,0,1,0,0,1,0,0,
???????????????????0,1,0,0,0,0,1,0,
???????????????????1,1,1,0,0,1,1,1,
???????????????????0,0,1,0,0,1,0,0,
???????????????????0,0,1,0,0,1,0,0,
???????????????????0,0,1,0,0,1,0,0,
???????????????????0,0,1,1,1,1,0,0,};
//輸出數組初始為-1,即該位置不顯示數字
????int buf_B[] = {-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1,
???????????????????-1,-1,-1,-1,-1,-1,-1,-1};
????//step 1:初始化OpenCL
????err = clGetPlatformIDs(1,&platform_id,NULL);
????if(err!=CL_SUCCESS)
????{
????????cout<<"clGetPlatformIDs error:"<<err<<endl;
????????return 0;
????}
????//這次我們只用CPU來進行并行運算，當然你也可以該成GPU
????clGetDeviceIDs(platform_id,CL_DEVICE_TYPE_GPU,1,&device,NULL);
????//step 2:創建上下文
????context = clCreateContext(NULL,1,&device,NULL,NULL,NULL);
????//step 3:創建命令隊列
????cmdQueue = clCreateCommandQueue(context,device,0,NULL);
????//step 4:創建數據緩沖區
????bufferA = clCreateBuffer(context,
?????????????????????????????CL_MEM_READ_ONLY,
?????????????????????????????datasize,NULL,NULL);
????bufferB = clCreateBuffer(context,
?????????????????????????????CL_MEM_WRITE_ONLY,
?????????????????????????????datasize,NULL,NULL);
????//step 5:將數據上傳到緩沖區
????clEnqueueWriteBuffer(cmdQueue,
?????????????????????????bufferA,CL_FALSE,
?????????????????????????0,datasize,
?????????????????????????buf_A,0,
?????????????????????????NULL,NULL);
????//step 6:加載編譯代碼,創建內核調用函數
????program = clCreateProgramWithSource(context,1,
????????????????????????????????????????(const char**)&buf_code,
????????????????????????????????????????NULL,NULL);
????clBuildProgram(program,1,&device,NULL,NULL,NULL);
????kernel = clCreateKernel(program,"rotation",NULL);
????//step 7:設置參數，執行內核
????float sinangle = sinf(angle);
????float cosangle = cosf(angle);
????clSetKernelArg(kernel,0,sizeof(cl_mem),&bufferA);
????clSetKernelArg(kernel,1,sizeof(cl_mem),&bufferB);
????clSetKernelArg(kernel,2,sizeof(cl_int),&dim_x);
????clSetKernelArg(kernel,3,sizeof(cl_int),&dim_y);
????clSetKernelArg(kernel,4,sizeof(cl_float),&sinangle);
????clSetKernelArg(kernel,5,sizeof(cl_float),&cosangle);
????//注意這里第三個參數已經改成2，表示二維數據。
????clEnqueueNDRangeKernel(cmdQueue,kernel,
???????????????????????????2,NULL,
???????????????????????????globalWorkSize,
???????????????????????????NULL,0,NULL,NULL);
????//step 8:取回計算結果
????clEnqueueReadBuffer(cmdQueue,bufferB,CL_TRUE,0,
????????????????????????datasize,buf_B,0,NULL,NULL);
????//輸出計算結果
????for(n=0;n<dim_x;n++)
????{
????????for(m=0;m<dim_y;m++)
????????{
????????????if(buf_A[m+dim_x*n]==0)
????????????????cout<<"? ";
????????????else
????????????????cout<< buf_A[m+dim_x*n] <<" ";
????????}
????????cout<<endl;
????}
????cout<<endl<<"====Rotate 45===="<<endl<<endl;
????for(n=0;n<dim_x;n++)
????{
????????for(m=0;m<dim_y;m++)
????????{
????????????if(buf_B[m+dim_x*n]<=0)
????????????????cout<<"? ";
????????????else
????????????????cout<< buf_B[m+dim_x*n] <<" ";
????????}
????????cout<<endl;
????}
????//釋放所有調用和內存
????clReleaseKernel(kernel);
????clReleaseProgram(program);
????clReleaseCommandQueue(cmdQueue);
????clReleaseMemObject(bufferA);
????clReleaseMemObject(bufferB);
????clReleaseContext(context);
????delete buf_code;
????return 0;
}