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
代码下载: