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

代码下载：

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