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, //block
   CL_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 20

int 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, C
	for (int i = 0; i < BUFFER_SIZE; i++) {
		A[i] = i;
		B[i] = i * 2;
		C[i] = 0;
	}

	try {
		std::vector<cl::Platform> platformList;

		// Pick platform
		cl::Platform::get(&platformList);

		// Pick first platform
		cl_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 context
		std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();

		// Create and program from source
		cl::Program::Sources sources(1, std::make_pair(kernelSourceCode, 0));
		cl::Program program(context, sources);

		// Build program
		program.build(devices);

		// Create buffer for A and copy host contents
		cl::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 contents
		cl::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 C
		cl::Buffer cBuffer = cl::Buffer(
			context,
			CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
			BUFFER_SIZE * sizeof(int),
			(void*)&C[0]);

		// Create kernel object
		cl::Kernel kernel(program, "vadd");

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

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

		// Do the work
		queue.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 memory
		err = 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;
}