OpenCL 學習step by step (6) 旋轉圖像

      在本教程中，我們學習用opencl進行簡單的圖像處理，對一個圖片進行旋轉。圖片讀入、保存等工作，我們使用開源的FreeImage，下載地址：http://freeimage.sourceforge.net/

      首先我們建立一個gFreeImage類，用來裝入圖像，該類主要調用FreeImage的函數，首先會初始化FreeImage庫，然后根據文件名猜測圖像文件格式，最終load圖像文件到變量FIBITMAP *bitmap中去。同時，我們還定義了2個緩沖

unsigned char *imageData;
unsigned char *imageData4;

用來存放圖像數據，之所以定義imageData4,是因為通常的圖片文件，比如jpg,bmp都是3個通道，沒有包括alpha通道，但是在gpu中處理數據時候，通常以vector4或者vector的形式進行，不能以vector3進行，所以我們裝入圖像后，會把imageData指向圖像數據，同時生成包括alpha通道的圖像數據imageData4。

     另外，我們還定義了一個函數LoadImageGrey，該函數用來裝入灰度圖，灰度圖一個像素用一個uchar表示。

在main.cpp中，我們首先定義一個cpu處理圖像旋轉的函數：

//CPU旋轉圖像
void cpu_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];
            }
        }
    }

    在main函數中，我們首先會裝入圖像文件，代碼如下：

int W, H;
gFreeImage img;
if(!img.LoadImageGrey("lenna.jpg"))
    {
    printf("can‘t load lenna.jpg\n");
    exit(0);
    }
else
    src_image = img.getImageDataGrey(W, H);

size_t mem_size = W*H;
cpu_image = (unsigned char*)malloc(mem_size);

    之后，定義2個cl memory對象，一個用來放原始圖像，一個用來放旋轉后的圖像。

//創建2個OpenCL內存對象
cl_mem d_ip = clCreateBuffer(
    context, CL_MEM_READ_ONLY,
    mem_size,
    NULL, NULL);
cl_mem d_op = clCreateBuffer(
    context, CL_MEM_WRITE_ONLY,
    mem_size,
    NULL, NULL);

cl_event writeEvt;
status = clEnqueueWriteBuffer (   
    queue , d_ip, CL_TRUE,
    0, mem_size, (void *)src_image,
    0, NULL, &writeEvt);
//等待數據傳輸完成再繼續往下執行
status = clFlush(queue);
waitForEventAndRelease(&writeEvt);
//clWaitForEvents(1, &writeEvt);

   旋轉kernel函數需要傳入6個參數：

//創建Kernel對象
cl_kernel kernel = clCreateKernel( program, "image_rotate", NULL );
//設置Kernel參數
float sintheta = 1, costheta = 0;
clSetKernelArg(kernel, 0, sizeof(cl_mem),  (void *)&d_ip);
clSetKernelArg(kernel, 1, sizeof(cl_mem),  (void *)&d_op);
clSetKernelArg(kernel, 2, sizeof(cl_int),  (void *)&W);
clSetKernelArg(kernel, 3, sizeof(cl_int),  (void *)&H);
clSetKernelArg(kernel, 4, sizeof(cl_float), (void *)&sintheta);
clSetKernelArg(kernel, 5, sizeof(cl_float), (void *)&costheta);
kernel執行的代碼為：

//執行kernel,Range用2維，work itmes size為W*H,
cl_event ev;
size_t globalThreads[] = {W, H};
size_t localThreads[] = {16, 16}; // localx*localy應該是64的倍數
printf("global_work_size =(%d,%d), local_work_size=(16, 16)\n",W,H);

clTimer.Reset();
clTimer.Start();
clEnqueueNDRangeKernel( queue,
    kernel,
    2,
    NULL,
    globalThreads,
    localThreads, 0, NULL, &ev);
//沒有設置local group size時候，系統將會自動設置為 (256,1)
status = clFlush( queue );
waitForEventAndRelease(&ev);
//clWaitForEvents(1, &ev);

clTimer.Stop();
printf("kernal total time:%.6f ms \n ", clTimer.GetElapsedTime()*1000 );

kernel函數代碼為：

#pragma OPENCL EXTENSION cl_amd_printf : enable
__kernel  void image_rotate(    __global uchar * src_data, __global uchar * dest_data,    //源圖像和輸出圖像都放在global memory中
                                                int W,    int H,                                //圖像size
                                               float sinTheta, float cosTheta )                    //旋轉角度
{    
    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]; 
     }
}

gpu執行完畢后，旋轉后的圖像保存在lenna_rotate.jpg,我們還會用cpu rotate函數執行一次旋轉，同時把生成的圖像保存到cpu_lenna_rotate.jpg。

完整的代碼請參考：

工程文件gclTutorial5

代碼下載：

http://files.cnblogs.com/mikewolf2002/gclTutorial.zip