OpenCL 教学(一)
Contents
- OpenCL简介
- OpenCL的架构
- OpenCL环境设定
- 开始撰写OpenCL程式
- 建立Command Queue
- 产生资料
- 配置记忆体并复制资料
- 编译OpenCL kernel程式
- 执行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 中,则大致的流程是:
- 把OpenCL 装置初始化。
- 在OpenCL 装置上配置三块记忆体,以存放a、b、c 三个阵列的资料。
- 把a 阵列和b 阵列的内容,复制到OpenCL 装置上。
- 编译要执行的OpenCL 程式(称为kernel)。
- 执行编译好的kernel。
- 把计算结果从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_a
和cl_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_base::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_base::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;
}
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 )
- cltut_1.zip - on Feb 3, 2010 8:54 AM by Chen Ping-Che (version 2 / earlier versions ) 7kDownload