[转]OpenCL 教学(一)

Contents

  1. OpenCL简介
  2. OpenCL的架构
  3. OpenCL环境设定
  4. 开始撰写OpenCL程式
  5. 建立Command Queue
  6. 产生资料
  7. 配置记忆体并复制资料
  8. 编译OpenCL kernel程式
  9. 执行OpenCL kernel
penCL 简介
OpenCL是由Khronos Group针对异质性计算装置(heterogeneous device)进行平行化运算所设计的标准API以及程式语言。所谓的「异质性计算装置」,是指在同一个电脑系统中,有两种以上架构差异很大的计算装置,例如一般的CPU以及显示晶片,或是像CELL的PPE以及SPE。目前,最为常见的就是所谓的GPGPU应用,也就是利用一般的显示晶片(即GPU)进行3D绘图以外的计算工作。

过去GPGPU的应用,有各种不同的使用方式。最早的GPGPU,多半是直接透过3D绘图的API进行,例如OpenGL或D3D的HLSL(High Level Shading Language)。但是,这样做有很多缺点,主要是即使想要进行的运算和3D绘图无关,仍需要处理很多3D绘图方面的动作(例如建立texture,安排render-to-texture动作等等)。这让GPGPU变得十分复杂。后来开始有些尝试把这些3D绘图部份隐藏起来的想法,例如由Stanford大学设计的BrookGPU,可以透过不同的backend将Brook程式转换成由CPU、Direct3D、或OpenGL来执行。另外,也有各家显示卡厂商自行开发的系统,包括ATI针对其产品设计的Close to Metal(以及后来的AMD Stream),以及NVIDIA的CUDA。Microsoft也在DirectX 11中加入了特别为GPGPU设计的DirectCompute。

由于各家厂商的GPGPU 方案都是互不相容的(例如AMD Stream 的程式无法在NVIDIA 的显示晶片上执行,而CUDA 的程式也不能在AMD 的显示晶片上执行),这对GPGPU 的发展是不利的,因为程式开发者必须为不同厂商的显示晶片分别撰写程式,或是选择只支援某个显示晶片厂商。由于显示晶片的发展愈来愈弹性化,GPGPU 的应用范围也增加,因此Apple 决定提出一个统一的GPGPU 方案。这个方案得到包括AMD、IBM、Intel、NVIDIA 等相关厂商的支持,并很快就交由Khronos Group 进行标准化。整个计画只花了五个月的时间,并在2008 年十二月时正式公开。第一个正式支援OpenCL 的作业系统是Apple 的MacOS X 10.6 "Snow Leopard"。AMD 和NVIDIA 也随后推出了在Windows 及Linux 上的OpenCL 实作。IBM 也推出了支援CELL 的OpenCL 实作。

OpenCL 的主要设计目的,是要提供一个容易使用、且适用于各种不同装置的平行化计算平台。因此,它提供了两种平行化的模式,包括task parallel 以及data parallel。目前GPGPU 的应用,主要是以data parallel 为主,这里也是以这个部份为主要重点。所谓的data parallel,指的是有大量的资料,都进行同样的处理。这种形式的平行化,在很多工作上都可以见到。例如,影像处理的程式,经常要对一个影像的每个pixel 进行同样的动作(例如Gaussian blur)。因此,这类工作很适合data parallel 的模式。

OpenCL 的架构

OpenCL 包括一组API 和一个程式语言。基本上,程式透过OpenCL API 取得OpenCL 装置(例如显示晶片)的相关资料,并将要在装置上执行的程式(使用OpenCL 程式语言撰写)编绎成适当的格式,在装置上执行。OpenCL API 也提供许多装置控制方面的动作,例如在OpenCL 装置上取得一块记忆体、把资料从主记忆体复制到OpenCL 装置上(或从OpenCL 装置上复制到主记忆体中)、取得装置动作的资讯(例如上一个程式执行所花费的时间)等等。

例如,我们先考虑一个简单的工作:把一群数字相加。在一般的C 程式中,可能是如下:

float a[DATA_SIZE];

float b[DATA_SIZE];

float result[DATA_SIZE];

// ...

for(int i = 0; i < DATA_SIZE; i++) {

result[i] = a[i] + b[i];

}
在OpenCL 中,则大致的流程是:
  1. 把OpenCL 装置初始化。
  2. 在OpenCL 装置上配置三块记忆体,以存放a、b、c 三个阵列的资料。
  3. 把a 阵列和b 阵列的内容,复制到OpenCL 装置上。
  4. 编译要执行的OpenCL 程式(称为kernel)。
  5. 执行编译好的kernel。
  6. 把计算结果从OpenCL 装置上,复制到result 阵列中。

透过data parallel 的模式,这里的OpenCL 程式非常简单,如下所示:

__kernel void adder(__global const float* a, __global const float* b, __global float* result)

{

int idx = get_global_id(0);

result[idx] = a[idx] + b[idx];

}

在一般的版本中,是透过一个回圈,执行DATA_SIZE次数的加法动作。而在OpenCL中,则是建立DATA_SIZE个work item,每个work item都执行上面所示的kernel。可以看到,OpenCL程式语言和一般的C语言非常类似。__kernel表示这个函式是在OpenCL装置上执行的。__global则表示这个指标是在global memory中(即OpenCL装置上的主要记忆体)。而get_global_id(0)会传回work item的编号,例如,如果有1024个work item,则编号会分别是0 ~ 1023(实际上编号可以是二维或三维,但在这里先只考虑一维的情形)。

要如何让上面这个简单的OpenCL kernel 实际在OpenCL 装置上执行呢?这就需要透过OpenCL API 的帮助了。以下会一步一步说明使用OpenCL API 的方法。

OpenCL 环境设定

在使用OpenCL API 之前,不免要进行一些环境的设定。相关的动作可以参考下列的文章:

开始撰写OpenCL 程式

在使用OpenCL API之前,和绝大部份所有其它的API一样,都需要include相关的header档案。由于在MacOS X 10.6下OpenCL的header档案命名方式和在其它作业系统下不同,因此,通常要使用一个#ifdef来进行区分。如下所示:

#ifdef __APPLE__

#include <OpenCL/opencl.h>

#else

#include <CL/cl.h>

#endif

这样就可以在MacOS X 10.6 下,以及其它的作业系统下,都可以include 正确的OpenCL header 档。

接着,要先取得系统上所有的OpenCL platform。在MacOS X 10.6 下,目前只有一个由Apple 提供的OpenCL platform,但是在其它系统上,可能会有不同厂商提供的多个不同的OpenCL platform,因此需要先取得platform 的数目:

cl_int err;

cl_uint num;

err = clGetPlatformIDs(0, 0, &num);

if(err != CL_SUCCESS) {

std::cerr << "Unable to get platforms\n";

return 0;

}

大部份的OpenCL API 会传回错误值。如果传回值是CL_SUCCESS 则表示执行成功,否则会传回某个错误值,表示失败的原因。

接着,再取得platform 的ID,这在建立OpenCL context 时会用到:

std::vector<cl_platform_id> platforms(num);

err = clGetPlatformIDs(num, &platforms[0], &num);

if(err != CL_SUCCESS) {

std::cerr << "Unable to get platform ID\n";

return 0;

}

在OpenCL 中,类似这样的模式很常出现:先呼叫第一次以取得数目,以便配置足够的记忆体量。接着,再呼叫第二次,取得实际的资料。

接下来,要建立一个OpenCL context。如下:

cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };

cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL);

if(context == 0) {

std::cerr << "Can't create OpenCL context\n";

return 0;

}
clReleaseContext(context);
return 0;
在上面的程式中,clCreateContextFromType是一个OpenCL的API,它可以从指定的装置类别中,建立一个OpenCL context。第一个参数是指定context的property。在OpenCL中,是透过一个property的阵列,以「property种类」及「property内容」成对出现,并以0做为结束。例如,以上面的例子来说,要指定的property种类是CL_CONTEXT_PLATFORM,即要使用的platform ID,而property内容则是由之前取得的platform ID中的第一个(即platforms[0])。由于property的内容可能是不同的资料型态,因此需要使用reinterpret_cast来进行强制转型。

第二个参数可以指定要使用的装置类别。目前可以使用的类别包括:

  • CL_DEVICE_TYPE_CPU:使用CPU 装置
  • CL_DEVICE_TYPE_GPU:使用显示晶片装置
  • CL_DEVICE_TYPE_ACCELERATOR:特定的OpenCL 加速装置,例如CELL
  • CL_DEVICE_TYPE_DEFAULT:系统预设的OpenCL 装置
  • CL_DEVICE_TYPE_ALL:所有系统中的OpenCL 装置

这里使用的是CL_DEVICE_TYPE_DEFAULT,也就是指定使用预设的装置。另外,在这里,直接使用了之前取得的OpenCL platform ID中的第一个ID(实际的程式中,可能会需要让使用者可以指定要使用哪一个platform)。

如果建立OpenCL context失败,会传回0。因此,要进行检查,并显示错误讯息。如果建立成功的话,在使用完后,要记得将context释放。这可以透过呼叫clReleaseContext来达成。

这个程式基本上已经可以编译执行了,但是当然它并没有真的做什么事情。

一个OpenCL context中可以包括一个或多个装置,所以接下来的工作是要取得装置的列表。要取得任何和OpenCL context相关的资料,可以使用clGetContextInfo函式。以下是取得装置列表的方式:

size_t cb;

clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb);

std::vector<cl_device_id> devices(cb / sizeof(cl_device_id));

clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0);

CL_CONTEXT_DEVICES表示要取得装置的列表。和前面取得platform ID的情形相同,clGetContextInfo被呼叫了两次:第一次是要取得需要存放装置列表所需的记忆体空间大小(也就是传入&cb),然后第二次呼叫才真正取得所有装置的列表。

接下来,可能会想要确定倒底找到的OpenCL装置是什么。所以,可以透过OpenCL API取得装置的名称,并将它印出来。取得和装置相关的资料,是使用clGetDeviceInfo函式,和前面的clGetContextInfo函式相当类似。以下是取得装置名称的方式:

clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb);

std::string devname;

devname.resize(cb);

clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0);

std::cout << "Device: " << devname.c_str() << "\n";

到目前为止,完整的程式应该如下所示:

// OpenCL tutorial 1

#include <iostream>

#include <string>

#include <vector>

#ifdef __APPLE__

#include <OpenCL/opencl.h>

#else

#include <CL/cl.h>

#endif

int main()

{

cl_int err;

cl_uint num;

err = clGetPlatformIDs(0, 0, &num);

if(err != CL_SUCCESS) {

std::cerr << "Unable to get platforms\n";

return 0;

}


std::vector<cl_platform_id> platforms(num);

err = clGetPlatformIDs(num, &platforms[0], &num);

if(err != CL_SUCCESS) {

std::cerr << "Unable to get platform ID\n";

return 0;

}


cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };

cl_context context = clCreateContextFromType( prop , CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL);
if(context == 0) {
std::cerr << "Can't create OpenCL context\n";
return 0;
}
size_t cb;
clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb);
std::vector<cl_device_id> devices(cb / sizeof(cl_device_id));
clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0);
clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb);
std::string devname;
devname.resize(cb);
clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0);
std::cout << "Device: " << devname.c_str() << "\n";
clReleaseContext(context);
return 0;

}

执行这个程式,如果建立OpenCL context 成功的话,应该会显示出找到的OpenCL 装置的名称,例如

Device: GeForce GTX 285

建立Command Queue

大部份OpenCL 的操作,都要透过command queue。Command queue 可以接收对一个OpenCL 装置的各种操作,并按照顺序执行(OpenCL 也容许把一个command queue 指定成不照顺序执行,即out-of-order execution,但是这里先不讨论这个使用方式)。所以,下一步是建立一个command queue:

cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, 0);

if(queue == 0) {

std::cerr << "Can't create command queue\n";
clReleaseContext(context);

return 0;

}
和context 一样,在程式结束前,要把command queue 释放,即:

clReleaseCommandQueue(queue);

上面的程式中,是把装置列表中的第一个装置(即devices[0])建立command queue。如果想要同时使用多个OpenCL装置,则每个装置都要有自己的command queue。
产生资料

由于这个程式的目的是要把一大堆数字进行相加,所以需要产生一些「测试资料」:

const int DATA_SIZE = 1048576;

std::vector<float> a(DATA_SIZE), b(DATA_SIZE), res(DATA_SIZE);

for(int i = 0; i < DATA_SIZE; i++) {

a[i] = std::rand();

b[i] = std::rand();

}
配置记忆体并复制资料

要使用OpenCL 装置进行运​​算时,通常会需要在OpenCL 装置上配置记忆体,并把资料从主记忆体中复制到装置上。有些OpenCL 装置可以直接从主记忆体存取资料,但是速度通常会比较慢,因为OpenCL 装置(例如显示卡)通常会有专用的高速记忆体。以下的程式配置三块记忆体:

cl_mem cl_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL);

cl_mem cl_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL);

cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);

if(cl_a == 0 || cl_b == 0 || cl_res == 0) {

std::cerr << "Can't create OpenCL buffer\n";

clReleaseMemObject(cl_a);

clReleaseMemObject(cl_b);

clReleaseMemObject(cl_res);

clReleaseCommandQueue(queue);

clReleaseContext(context);

return 0;

}
clCreateBuffer函式可以用来配置记忆体。它的第二个参数可以指定记忆体的使用方式,包括:
  • CL_MEM_READ_ONLY:表示OpenCL kernel 只会对这块记忆体进行读取的动作
  • CL_MEM_WRITE_ONLY:表示OpenCL kernel 只会对这块记忆体进行写入的动作
  • CL_MEM_READ_WRITE:表示OpenCL kernel 会对这块记忆体进行读取和写入的动作
  • CL_MEM_USE_HOST_PTR:表示希望OpenCL 装置直接使用指定的主记忆体位址。要注意的是,如果OpenCL 装置无法直接存取主记忆体,它可能会将指定的主记忆体位址的资料复制到OpenCL 装置上。
  • CL_MEM_ALLOC_HOST_PTR:表示希望配置的记忆体是在主记忆体中,而不是在OpenCL 装置上。不能和CL_MEM_USE_HOST_PTR 同时使用。
  • CL_MEM_COPY_HOST_PTR:将指定的主记忆体位址的资料,复制到配置好的记忆体中。不能和CL_MEM_USE_HOST_PTR 同时使用。

第三个参数是指定要配置的记忆体大小,以bytes为单位。在上面的程式中,指定的大小是sizeof(cl_float) * DATA_SIZE

第四个参数是指定主记忆体的位置。因为对cl_acl_b来说,在第二个参数中,指定了CL_MEM_COPY_HOST_PTR,因此要指定想要复制的资料的位址。cl_res则不需要指定。

第五个参数是指定错误码的传回位址。在这里并没有使用到。

如果clCreateBuffer因为某些原因无法配置记忆体(例如OpenCL装置上的记忆体不够),则会传回0。要释放配置的记忆体,可以使用clReleaseMemObject函式。

编译OpenCL kernel 程式

现在执行OpenCL kernel 的准备工作已经大致完成了。所以,现在剩下的工作,就是把OpenCL kernel 程式编释并执行。首先,先把前面提过的OpenCL kernel 程式,存放在一个文字档中,命名为shader.cl:

__kernel void adder(__global const float* a, __global const float* b, __global float* result)

{

int idx = get_global_id(0);

result[idx] = a[idx] + b[idx];

}

要编译这个kernel程式,首先要把档案内容读进来,再使用clCreateProgramWithSource这个函式,然后再使用clBuildProgram编译。如下所示:

cl_program load_program(cl_context context, const char* filename)

{

std::ifstream in(filename, std::ios_bas​​e::binary);

if(!in.good()) {

return 0;

}
// get file length
in.seekg(0, std::ios_base::end);
size_t length = in.tellg();
in.seekg(0, std::ios_base::beg);
// read program source
std::vector<char> data(length + 1);
in.read(&data[0], length);
data[length] = 0;
// create and build program 
const char* source = &data[0];
cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0);
if(program == 0) {
return 0;
}
if(clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS) {
return 0;
}
return program;
}
上面的程式,就是直接将档案读到记忆体中,再呼叫clCreateProgramWithSource建立一个program object。建立成功后,再呼叫clBuildProgram函式编译程式。clBuildProgram函式可以指定很多参数,不过在这里暂时没有使用到。

有了这个函式,在main 函式中,直接呼叫:

cl_program program = load_program(context, "shader.cl");

if(program == 0) {

std::cerr << "Can't load or build program\n";

clReleaseMemObject(cl_a);

clReleaseMemObject(cl_b);

clReleaseMemObject(cl_res);

clReleaseCommandQueue(queue);

clReleaseContext(context);

return 0;

}

同样的,在程式结束前,要记得将program object 释放:

clReleaseProgram(program);

一个OpenCL kernel 程式里面可以有很多个函式。因此,还要取得程式中函式的进入点:

cl_kernel adder = clCreateKernel(program, "adder", 0);

if(adder == 0) {

std::cerr << "Can't load kernel\n";

clReleaseProgram(program);

clReleaseMemObject(cl_a);

clReleaseMemObject(cl_b);

clReleaseMemObject(cl_res);

clReleaseCommandQueue(queue);

clReleaseContext(context);

return 0;

}
和program object 一样,取得的kernel object 也需要在程式结束前释放:

clReleaseKernel(adder);

执行OpenCL kernel

弄了这么多,总算可以执行OpenCL kernel程式了。要执行kernel程式,只需要先设定好函式的参数。adder函式有三个参数要设定:

clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a);

clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b);

clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res);

设定参数是使用clSetKernelArg函式。它的参数很简单:第一个参数是要设定的kernel object,第二个是参数的编号(从0开始),第三个参数是要设定的参数的大小,第四个参数则是实际上要设定的参数内部。以这里的adder函式来说,三个参数都是指向memory object的指标。

设定好参数后,就可以开始执行了。如下:

size_t work_size = DATA_SIZE;

err = clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, 0);

clEnqueueNDRangeKernel会把执行一个kernel的动作加到command queue里面。第三个参数(1)是指定work item数目的维度,在这里就是一维。第五个参数是指定work item的总数目,也就是DATA_SIZE。后面的参数现在暂时先不用管。如果成功加入的话,会传回CL_SUCCESS。否则会传回错误值。

在执行kernel 被加到command queue 之后,就可能会开始执行(如果command queue 现在没有别的工作的话)。但是clEnqueueNDRangeKernel 是非同步的,也就是说,它并不会等待OpenCL 装置执行完毕才传回。这样可以让CPU 在OpenCL 装置在进行运算的同时,进行其它的动作。

由于执行的结果是在OpenCL 装置的记忆体中,所以要取得结果,需要把它的内容复制到CPU 能存取的主记忆体中。这可以透过下面的程式完成:

if(err == CL_SUCCESS) {

err = clEnqueueReadBuffer(queue, cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0);

}
clEnqueueReadBuffer函式会把「将记忆体资料从OpenCL装置复制到主记忆体」的动作加到command queue中。第三个参数表示是否要等待复制的动作完成才传回,CL_TRUE表示要等待。第五个参数是要复制的资料大小,第六个参数则是目标的位址。

由于这里指定要等待复制动作完成,所以当函式传回时,资料已经完全复制完成了。最后是进行验证,确定资料正确:

if(err == CL_SUCCESS) {

bool correct = true;

for(int i = 0; i < DATA_SIZE; i++) {

if(a[i] + b[i] != res[i]) {

correct = false;

break;

}
}
if(correct) {
std::cout << "Data is correct\n";
}
else {
std::cout << "Data is incorrect\n";
}
}
else {
std::cerr << "Can't run kernel or read back data\n";
}
到这里,整个程式就算是完成了。编译后执行,如果顺利的话,应该会印出

Data is correct

的讯息。

以下是整个程式的全貌:

// OpenCL tutorial 1

#include <iostream>

#include <fstream>

#include <string>

#include <vector>

#include <cstdlib>

#ifdef __APPLE__

#include <OpenCL/opencl.h>

#else

#include <CL/cl.h>

#endif

cl_program load_program(cl_context context, const char* filename)

{

std::ifstream in(filename, std::ios_bas​​e::binary);
if(!in.good()) {
return 0;
}
// get file length
in.seekg(0, std::ios_bas​​e::end);
size_t len​​gth = in.tellg();
in.seekg(0, std::ios_bas​​e::beg);
// read program source
std::vector<char> data(length + 1);
in.read(&data[0], length);
data[length] = 0;
// create and build program
const char* source = &data[0];
cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0);
if(program == 0) {
return 0;
}
if(clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS) {
return 0;
}
return program;

}

int main()

{

cl_int err;

cl_uint num;

err = clGetPlatformIDs(0, 0, &num);

if(err != CL_SUCCESS) {

std::cerr << "Unable to get platforms\n";

return 0;

}


std::vector<cl_platform_id> platforms(num);

err = clGetPlatformIDs(num, &platforms[0], &num);

if(err != CL_SUCCESS) {

std::cerr << "Unable to get platform ID\n";

return 0;

}


cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };

cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL);
if(context == 0) {
std::cerr << "Can't create OpenCL context\n";
return 0;
}
size_t cb;
clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb);
std::vector<cl_device_id> devices(cb / sizeof(cl_device_id));
clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0);
clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb);
std::string devname;
devname.resize(cb);
clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0);
std::cout << "Device: " << devname.c_str() << "\n";
cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, 0);
if(queue == 0) {
std::cerr << "Can't create command queue\n";
clReleaseContext(context);
return 0;
}
const int DATA_SIZE = 1048576;
std::vector<float> a(DATA_SIZE), b(DATA_SIZE), res(DATA_SIZE);
for(int i = 0; i < DATA_SIZE; i++) {
a[i] = std::rand();
b[i] = std::rand();
}
cl_mem cl_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL);
cl_mem cl_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL);
cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);
if(cl_a == 0 || cl_b == 0 || cl_res == 0) {
std::cerr << "Can't create OpenCL buffer\n";
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
cl_program program = load_program(context, "shader.cl");
if(program == 0) {
std::cerr << "Can't load or build program\n";
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
cl_kernel adder = clCreateKernel(program, "adder", 0);
if(adder == 0) {
std::cerr << "Can't load kernel\n";
clReleaseProgram(program);
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a);
clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b);
clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res);
size_t work_size = DATA_SIZE;
err = clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, 0);
if(err == CL_SUCCESS) {
err = clEnqueueReadBuffer(queue, cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0);
}
if(err == CL_SUCCESS) {
bool correct = true;
for(int i = 0; i < DATA_SIZE; i++) {
if(a[i] + b[i] != res[i]) {
correct = false;
break;
}
}
if(correct) {
std::cout << "Data is correct\n";
}
else {
std::cout << "Data is incorrect\n";
}
}
else {
std::cerr << "Can't run kernel or read back data\n";
}
clReleaseKernel(adder);
clReleaseProgram(program);
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;

}

在附件中可以下载包括Xcode project 以及Visual Studio 2008 project 档的原始码。

Attachments ( 1 )

登录 最近的站点活动 服务条款 举报不良信息 打印页面 Google协作平台强力驱动

posted @ 2011-10-19 23:44  Let it be!  阅读(4481)  评论(0编辑  收藏  举报