OpenCL编程指南-10.2使用C++包装器API的矢量相加示例

选择OpenCL平台并创建一个上下文

建立OpenCL的第一步是选择一个平台。第2章介绍过，OpenCL使用了ICD模型，其中可以有多个OpenCL实现在一个系统上并存。类似于HelloWorld示例，这个矢量相加程序展示了选择OpenCL平台的一种最简单的方法：选择第一个可用的平台。

首先，调用cl::Platform::get()得到平台列表：

std::vector<cl::Platform> platformList;
cl::Platform::get(&platformList);

得到平台列表之后，这个例子会调用cl::Context()创建一个上下文。cl::Context ()调用会尝试由一个GPU设备创建上下文。如果失败，程序会产生一个异常，这个程序使用了OpenCL C++包装器异常特性，以一个错误消息终止。创建上下文的代码如下：

cl_context_properties cprops[] = (CL_CONTEXT_PLATFORM,(cl_context_properties)(platformList[0])(),0);
cl::Context context(CL_DEVICE_TYPE_GPU, cprops);   

选择一个设备并创建命令队列

选择一个平台并创建上下文之后，矢量相加应用程序的下一步是选择一个设备，并创建一个命令队列。第一个任务是查询与之前所创建上下文关联的设备集合。可以通过cl::Context::getInfo<CL_CONTEXT_DEVICES >()调用来查询，这会返回与上下文关联的设备std::vector。

在继续学习后面的内容之前，先来了解getInfo()方法，因为它遵循了C++包装器API中通用的一种模式。一般来说，对于一个支持查询接口的CAPl对象（例如，查询接口为clGetXXInfo()，其中xx是所查询C API对象的名），任何表示这样一个CAPI对象的C++包装器API对象都有相应的一个接口，形式如下：

template <cl_int> typename
detail::param_traits<detail::cl_XX_info, name>::param_type
cl::Object::getInfo(void);

乍一看可能会让你有些害怕，因为这里使用了一种称为特征类(traits)的C++模板技术（这里用于关联clGetXXInfo()提供的共享功能），不过，由于使用这些getInfo()函数的程序在实际中从来不需要引用特征类组件，所以对于开发人员编写的代码没有任何影响。需要指出的重要一点是，所有对应一个底层C API对象的C++包装器API对象都有一个模板方法，名为getInfo()，以查询的cl_xx_info枚举值作为其模板参数。其效果是可以静态检查所请求的值是否合法，也就是说，一个特定的getInfo()方法只接受相应cl_xx_info枚举中定义的值。通过使用这种特征类技术，getInfo()函数可以自动推导出结果类型。

再来看矢量相加示例，要为关联的一组设备查询一个上下文，可以用CL_CONTEXT_DEVICES限定相应的cl::Context::getInfo()，返回std::vector<cl::Device>。通过以下代码可以说明：

//Query the set of devices attached to the context
std::vector<cl::Device> device =context.getInfo<CL_CONTEXT_DEVICES>();

注意：利用C++包装器API查询方法，现在不再需要先查询上下文来找出需要多大的空间存储设备列表，然后再提供另一个查询调用得到具体的设备。所有这些都隐藏在C++包装器API的一个简单的通用接口中。

选择设备集合之后，可以用cl::CommandQueue()创建一个命令队列，为简单起见，这里选择第一个设备：

//Create command-queue
cl::CommandQueue queue(context, device[0], 0);

创建和构建程序对象

矢量相加示例中的下一步是使用cl::Program()由OpenCL C内核源代码创建一个程序对象（矢量相加示例的内核源代码在本章最后的代码清单2-1中给出，这里不再重复）。程序对象用内核源代码加载，然后使用cl::Program::build()编译这个代码，以便在与上下文关联的设备上执行。下面给出相应的代码：

cl::Program::Sources sources(1,std::make_pair(kernelSourceCode,0));
cl::Program program(context, sources);
program.build(devices);

与其他C++包装器API调用类似，如果出现错误，则会有一个异常，程序将退出。

创建内核和内存对象

要执行OpenCL计算内核，需要在OpenCL设备上可访问的内存中分配内核函数的参数，这里就是缓冲区对象。这些缓冲区对象使用cl::Buffer()创建。对于输入缓冲区，我们使用CL_MEM_COPY_FROM_HOST_PTR来避免额外的调用来移动输人数据。对于输出缓冲区（即矢量相加的结果），则使用CL_MEM_USE_HOST_PTR，这要求将结果缓冲区映射到宿主机内存以便访问结果。可以使用以下代码来分配这些缓冲区：

cl::Buffer aBuffer = cl::Buffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,BUFFER_SIZE * sizeof(int),(void *) &A[0]);cl::Buffer bBuffer = cl::Buffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,BUFFER_SIZE * sizeof(int),(void *) &B[0]);cl::Buffer cBuffer = cl::Buffer(context,CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,BUFFER_SIZE * sizeof(int),(void *) &C[0]);

利用cl::Kernel()调用创建内核对象：

cl::Kernel kernel(program, "vadd");

执行矢量相加内核

既然已经创建了内核和内存对象，矢量相加程序终于可以将内核人队等待执行了。内核函数的所有参数都要使用cl::Kernel:setArg()方法来设置。根据C API中的clSetKernelArg()，这个函数的第一个参数是内核函数参数索引。vadd()内核有3个参数（a、b和c)，分别对应索引0、1和2。将之前创建的内存对象传入这个内核对象：

kernel.setArg(0, aBuffer);
kernel.setArg(1, bBuffer);
kernel.setArg(2, cBuffer);

与以往一样，设置内核参数之后，矢量相加示例使用命令队列将内核入队等待在设备上执行。这是通过调用cl::CommandQueue::enqueueNDRangeKernel()完成的。全局和局部工作大小使用cl::Range()传递。

对于局部工作大小，使用cl::Range()对象的一个特殊实例cl::NullRange，顾名思义，它对应于C API中传递NULL，允许运行时为设备确定最佳的工作组大小和请求的全局工作大小。

queue.enqueueNDRangeKernel(kernel,cl::NullRange,cl::NDRange(BUFFER_SIZE),cl::NullRange);

将内核入队等待执行并不意味着内核会立即执行。可以使用cl::CommandQueue::flush()或cl::CommandQueue::finish()强制提交到设备立即执行。不过，由于这个矢量相加示例只是要显示结果，所以它使用了一个阻塞的cl::CommandQueue::enqueueMapBuffer()，将输出缓冲区映射到一个宿主机指针：

int * output = (int *)queue.enqueueMapBuffere(cBuffer,CL_TRUE, //blockCL_MAP_READ,0,BUFFER_SIZE * sizeof(int));

宿主机应用程序再处理output 指向的数据，一旦完成，必须用cl::CommandQueue::enqueueUnmapMemObj()调用释放映射的内存：

err = queue.enqueueUnmapMemObject(cBuffer,(void *)output);

代码示例

#define __CL_ENABLE_EXCEPTIONS#include <CL/cl.hpp>#include <cstdio>
#include <cstdlib>
#include <iostream>#define BUFFER_SIZE 20int A[BUFFER_SIZE];
int B[BUFFER_SIZE];
int C[BUFFER_SIZE];static char
kernelSourceCode[] =
"__kernel void                                                               \n"
"vadd(__global int * a, __global int * b, __global int * c)                                                                     \n"
"{                                                                           \n"
"    size_t i =  get_global_id(0);                                           \n"
"                                                                            \n"
"    c[i] = a[i] + b[i];                                                     \n"
"}                                                                           \n"
;int
main(void)
{cl_int err;// Initialize A, B, Cfor (int i = 0; i < BUFFER_SIZE; i++) {A[i] = i;B[i] = i * 2;C[i] = 0;}try {std::vector<cl::Platform> platformList;// Pick platformcl::Platform::get(&platformList);// Pick first platformcl_context_properties cprops[] = {CL_CONTEXT_PLATFORM, (cl_context_properties)(platformList[0])(), 0 };cl::Context context(CL_DEVICE_TYPE_GPU, cprops);// Query the set of devices attched to the contextstd::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();// Create and program from sourcecl::Program::Sources sources(1, std::make_pair(kernelSourceCode, 0));cl::Program program(context, sources);// Build programprogram.build(devices);// Create buffer for A and copy host contentscl::Buffer aBuffer = cl::Buffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,BUFFER_SIZE * sizeof(int),(void*)&A[0]);// Create buffer for B and copy host contentscl::Buffer bBuffer = cl::Buffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,BUFFER_SIZE * sizeof(int),(void*)&B[0]);// Create buffer for that uses the host ptr Ccl::Buffer cBuffer = cl::Buffer(context,CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,BUFFER_SIZE * sizeof(int),(void*)&C[0]);// Create kernel objectcl::Kernel kernel(program, "vadd");// Set kernel argskernel.setArg(0, aBuffer);kernel.setArg(1, bBuffer);kernel.setArg(2, cBuffer);// Create command queuecl::CommandQueue queue(context, devices[0], 0);// Do the workqueue.enqueueNDRangeKernel(kernel,cl::NullRange,cl::NDRange(BUFFER_SIZE),cl::NullRange);// Map cBuffer to host pointer. This enforces a sync with // the host backing space, remember we choose GPU device.int* output = (int*)queue.enqueueMapBuffer(cBuffer,CL_TRUE, // block CL_MAP_READ,0,BUFFER_SIZE * sizeof(int));for (int i = 0; i < BUFFER_SIZE; i++) {std::cout << C[i] << " ";}std::cout << std::endl;// Finally release our hold on accessing the memoryerr = queue.enqueueUnmapMemObject(cBuffer,(void*)output);// There is no need to perform a finish on the final unmap// or release any objects as this all happens implicitly with// the C++ Wrapper API.}catch (cl::Error err) {std::cerr<< "ERROR: "<< err.what()<< "("<< err.err()<< ")"<< std::endl;return EXIT_FAILURE;}return EXIT_SUCCESS;
}