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