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 之前，不免要进行一些环境的设定。相关的动作可以参考下列的文章：

• 在Windows 下使用OpenCL
• 在Xcode 中使用OpenCL

开始撰写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_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 )

• • cltut_1.zip - on Feb 3, 2010 8:54 AM by Chen Ping-Che (version 2 / earlier versions ) 7kDownload