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

      首先我們建立一個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函數中，我們首先會裝入圖像檔案，代碼如下：

    之後，定義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,

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函數代碼為：

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

完整的代碼請參考：

工程檔案gclTutorial5

代碼下載下傳：

<a href="http://files.cnblogs.com/mikewolf2002/gclTutorial.zip">http://files.cnblogs.com/mikewolf2002/gclTutorial.zip</a>