现在，我们开始写一个简单的opencl程序，计算两个数组相加的和，放到另一个数组中去。程序用cpu和gpu分别计算，最后验证它们是否相等。opencl程序的流程大致如下： 下面是source code中的主要代码： int main(int argc, char* argv[]) { //在host内存中创建
现在，我们开始写一个简单的opencl程序，计算两个数组相加的和，放到另一个数组中去。程序用cpu和gpu分别计算，最后验证它们是否相等。opencl程序的流程大致如下：
下面是source code中的主要代码：
int main(int argc, char* argv[])
{
//在host内存中创建三个缓冲区
float *buf1 = 0;
float *buf2 = 0;
float *buf = 0;
buf1 =(float *)malloc(bufsize * sizeof(float));
buf2 =(float *)malloc(bufsize * sizeof(float));
buf =(float *)malloc(bufsize * sizeof(float));
//用一些随机值初始化buf1和buf2的内容
int i;
srand( (unsigned)time( null ) );
for(i = 0; i
buf1[i] = rand()%65535;
srand( (unsigned)time( null ) +1000);
for(i = 0; i
buf2[i] = rand()%65535;
//cpu计算buf1,buf2的和
for(i = 0; i
buf[i] = buf1[i] + buf2[i];
cl_uint status;
cl_platform_id platform;
//创建平台对象
status = clgetplatformids( 1, &platform, null );
注意：如果我们系统中安装不止一个opencl平台，比如我的os中，有intel和amd两家opencl平台，用上面这行代码，有可能会出错，因为它得到了intel的opencl平台，而intel的平台只支持cpu，而我们后面的操作都是基于gpu，这时我们可以用下面的代码，得到amd的opencl平台。
cl_uint numplatforms;std::string platformvendor;
status = clgetplatformids(0, null, &numplatforms);
if(status != cl_success)
{
return 0;
}
if (0
{
cl_platform_id* platforms = new cl_platform_id[numplatforms];
status = clgetplatformids(numplatforms, platforms, null);
char platformname[100];
for (unsigned i = 0; i
{
status = clgetplatforminfo(platforms[i],
cl_platform_vendor,
sizeof(platformname),
platformname,
null);
platform = platforms[i];
platformvendor.assign(platformname);
if (!strcmp(platformname, advanced micro devices, inc.))
{
break;
}
}
std::cout platform found : \n;
delete[] platforms;
}
cl_device_id device;
//创建gpu设备
clgetdeviceids( platform, cl_device_type_gpu, 1, &device, null);
//创建context
cl_context context = clcreatecontext( null, 1, &device, null, null, null);
//创建命令队列
cl_command_queue queue = clcreatecommandqueue( context,
device,
cl_queue_profiling_enable, null );
//创建三个opencl内存对象，并把buf1的内容通过隐式拷贝的方式
//拷贝到clbuf1,buf2的内容通过显示拷贝的方式拷贝到clbuf2
cl_mem clbuf1 = clcreatebuffer(context,
cl_mem_read_only | cl_mem_copy_host_ptr,
bufsize*sizeof(cl_float),buf1,
null );
cl_mem clbuf2 = clcreatebuffer(context,
cl_mem_read_only ,
bufsize*sizeof(cl_float),null,
null );
cl_event writeevt;
status = clenqueuewritebuffer(queue, clbuf2, 1, 0, bufsize*sizeof(cl_float), buf2, 0, 0, 0);
上面这行代码把buf2中的内容拷贝到clbuf2,因为buf2位于host端，clbuf2位于device端，所以这个函数会执行一次host到device的传输操作，或者说一次system memory到video memory的拷贝操作，所以我在该函数的后面放置了clfush函数，表示把command queue中的所有命令提交到device(注意：该命令并不保证命令执行完成),所以我们调用函数waitforeventandrelease来等待write缓冲的完成，swaitforeventandreleae 是一个用户定义的函数，它的内容如下，主要代码就是通过event来查询我们的操作是否完成，没完成的话，程序就一直block在这行代码处，另外我们也可以用opencl中内置的函数clwaitforevents来代替clflush和swaitforeventandreleae。
//等待事件完成int waitforeventandrelease(cl_event *event)
{
cl_int status = cl_success;
cl_int eventstatus = cl_queued;
while(eventstatus != cl_complete)
{
status = clgeteventinfo(
*event,
cl_event_command_execution_status,
sizeof(cl_int),
&eventstatus,
null);
}
status = clreleaseevent(*event);
return 0;
}
status = clflush(queue);
//等待数据传输完成再继续往下执行
waitforeventandrelease(&writeevt);
cl_mem buffer = clcreatebuffer( context,
cl_mem_write_only,
bufsize * sizeof(cl_float),
null, null );
kernel文件中放的是gpu中执行的代码，它被放在一个单独的文件add.cl中，本程序中kernel代码非常简单，只是执行两个数组相加。kernel的代码为：
__kernel void vecadd(__global const float* a, __global const float* b, __global float* c){
int id = get_global_id(0);
c[id] = a[id] + b[id];
}
//kernel文件为add.cl
const char * filename = add.cl
std::string sourcestr;
status = converttostring(filename, sourcestr);
converttostring也是用户定义的函数，该函数把kernel源文件读入到一个string中，它的代码如下：
//把文本文件读入一个string中,用来读入kernel源文件int converttostring(const char *filename, std::string& s)
{
size_t size;
char* str;
std::fstream f(filename, (std::fstream::in | std::fstream::binary));
if(f.is_open())
{
size_t filesize;
f.seekg(0, std::fstream::end);
size = filesize = (size_t)f.tellg();
f.seekg(0, std::fstream::beg);
str = new char[size+1];
if(!str)
{
f.close();
return null;
}
f.read(str, filesize);
f.close();
str[size] = '\0';
s = str;
delete[] str;
return 0;
}
printf(error: failed to open file %s\n, filename);
return 1;
}
const char * source = sourcestr.c_str();
size_t sourcesize[] = { strlen(source) };
//创建程序对象
cl_program program = clcreateprogramwithsource( context, 1, &source, sourcesize, null);
//编译程序对象
status = clbuildprogram( program, 1, &device, null, null, null );
if(status != 0)
{
printf(clbuild failed:%d\n, status);
char tbuf[0x10000];
clgetprogrambuildinfo(program, device, cl_program_build_log, 0x10000, tbuf, null);
printf(\n%s\n, tbuf);
return -1;
}
//创建kernel对象
cl_kernel kernel = clcreatekernel( program, vecadd, null );
//设置kernel参数
cl_int clnum = bufsize;
clsetkernelarg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);
clsetkernelarg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);
clsetkernelarg(kernel, 2, sizeof(cl_mem), (void*) &buffer);
注意：在执行kernel时候，我们只设置了global work items数量，没有设置group size，这时候，系统会使用默认的work group size，通常可能是256之类的。
//执行kernel,range用1维，work itmes size为bufsize
cl_event ev;
size_t global_work_size = bufsize;
clenqueuendrangekernel( queue, kernel, 1, null, &global_work_size, null, 0, null, &ev);
status = clflush( queue );
waitforeventandrelease(&ev);
//数据拷回host内存
cl_float *ptr;
cl_event mapevt;
ptr = (cl_float *) clenqueuemapbuffer( queue, buffer, cl_true, cl_map_read, 0, bufsize * sizeof(cl_float), 0, null, null, null );
status = clflush( queue );
waitforeventandrelease(&mapevt);
//结果验证，和cpu计算的结果比较
if(!memcmp(buf, ptr, bufsize))
printf(verify passed\n);
else printf(verify failed);
if(buf)
free(buf);
if(buf1)
free(buf1);
if(buf2)
free(buf2);
程序结束后，这些opencl对象一般会自动释放，但是为了程序完整，养成一个好习惯，这儿我加上了手动释放opencl对象的代码。
//删除opencl资源对象
clreleasememobject(clbuf1);
clreleasememobject(clbuf2);
clreleasememobject(buffer);
clreleaseprogram(program);
clreleasecommandqueue(queue);
clreleasecontext(context);
return 0;
}
程序执行后的界面如下:
完整的代码请参考：
工程文件gcltutorial1
代码下载：http://files.cnblogs.com/mikewolf2002/gcltutorial.zip
原文作者：迈克老狼