OpenCL专题03：程序框架详解（以并行化向量相加为例）

最新推荐文章于 2022-08-18 20:05:13 发布

吴智深

最新推荐文章于 2022-08-18 20:05:13 发布

阅读量1.1k

点赞数 2

分类专栏： OpenCL 文章标签： c++ 算法开发语言

本文链接：https://blog.csdn.net/deepbodhi/article/details/125311685

版权

OpenCL 专栏收录该内容

4 篇文章 9 订阅

订阅专栏

本文详细介绍了如何使用OpenCL开发程序，以实现两个向量相加为例，涵盖了从获取Platform、Device，创建Context、CommandQueue，创建Buffer，到创建OpenCL程序、编译、创建kernel，执行kernel并读取结果的完整流程。通过这个案例，读者可以了解OpenCL编程的基本步骤和关键函数的使用。

摘要由CSDN通过智能技术生成

1. 基本框架

开发一个OpenCL程序，包括三大阶共10个步骤。
在这里插入图片描述

2. 分步详解

我们通过一个案例来对各步骤进行说明。案例是实现两个向量相加，所得结果放在第三个向量中。通过OpenCL，利用GPU并发执行。

2.0 引入头文件并定义一个异常类

#define CL_TARGET_OPENCL_VERSION 220
#include <stdio.h>
#include <string.h>
#include <stdlib.h>
#include <iostream>
#include <CL/cl.h>
using namespace std;
class OCLException
{
    public:
    int type;//0：Error, 1: Warning
    int ID;
    string content;
    OCLException(int type, int ID, string content)
    {
        this->type=type;
        this->ID=ID;
        this->content=content;
    }
};
int main()
{
	//程序主体部分，下文分步解释
}

这里定义了一个异常类OLException，统一处理程序运行中出现的错误、警告等。
该类中预设三个属性：

type：异常类型（0为错误，1为警告）
ID：异常编号（在编写程序时，同步维护异常码表）
content：异常说明

2.1 获取Platform

	cl_int ret;

    //step1: 获取平台
    cl_uint num_platforms;
    cl_platform_id platform_ids[5];
    char param_value[100];
    size_t param_value_size_ret;

    ret = clGetPlatformIDs(0, NULL, &num_platforms);
    if(ret!=CL_SUCCESS) throw new OCLException(0,1,"clGetPlatformsIDs Error");
    if(num_platforms==0) throw new OCLException(1,1,"no platform found");
    cout<<"Platform Number："<<num_platforms<<endl;

    ret = clGetPlatformIDs(num_platforms,platform_ids,NULL);
    if(ret!=CL_SUCCESS) throw new OCLException(0,2,"clGetPlatformsIDs Error");
    for(int i=0;i<num_platforms;i++)
    {
        cout<<"Platform ID:"<<platform_ids[i]<<endl;
        ret = clGetPlatformInfo(platform_ids[i],CL_PLATFORM_NAME,100,param_value,&param_value_size_ret);
        if(ret!=CL_SUCCESS) throw new OCLException(0,3,"clGetPlatformInfo Error");
        cout<<"Platform Name:"<<param_value<<endl;
    }

其中主要有两个函数clGetPlatformIDs和clGetPlatformInfo。

cl_int clGetPlatformIDs(
	cl_uint num_entries, 
	cl_platform_id* platforms, 
	cl_uint* num_platforms)

第一次调用该函数时，设置num_entries为0，platforms设置为NULL，函数将通过num_platforms返回系统中可用Platform的总数。
第二次调用该函数时，设置num_entries为刚才获取到的可用Platform总数，函数通过platforms返回Platform ID的数组。

cl_int clGetPlatformInfo(
	cl_platform_id platform, 
	cl_platform_info param_name, 
	size_t param_value_size, 
	void* param_value,
	size_t* param_value_size_ret)

platform设置为上一个函数获取到的Platform ID数组中的项
param_name表示需要查询的具体信息，文档中给出了8种类型，一般查询CL_PLATFORM_NAME就可以了
param_value_size指明用于接受查询信息的数组长度
param_value则是接受查询信息的数组
param_value_size_ret是实际返回信息的长度

2.2 获取Device

	//step2: 获取Device（以第1个Platform为例）
    cout<<"The 1st Platform:"<<endl;
    cl_uint num_devices;
    cl_device_id device_ids[5];
    char device_param_value[100];
    size_t device_param_value_size_ret;

    ret = clGetDeviceIDs(platform_ids[0],CL_DEVICE_TYPE_GPU,0,NULL,&num_devices);
    if(ret!=CL_SUCCESS) throw new OCLException(0,4,"clGetDeviceIDs Error");
    if(num_devices==0) throw new OCLException(1,2,"no device found");
    cout<<"Device Number:"<<num_devices<<endl;

    ret=clGetDeviceIDs(platform_ids[0],CL_DEVICE_TYPE_GPU,num_devices,device_ids,NULL);
    if(ret!=CL_SUCCESS) throw new OCLException(0,5,"clGetDeviceIDs Error");
    for(int i=0;i<num_devices;i++)
    {
        cout<<"Device ID:"<<device_ids[i]<<endl;
        ret = clGetDeviceInfo(device_ids[i],CL_DEVICE_NAME,100,device_param_value,&device_param_value_size_ret);
        if(ret!=CL_SUCCESS) throw new OCLException(0,6,"clGetDeviceInfo Error");
        cout<<"Device Name:"<<device_param_value<<endl;
    }

其中主要有两个函数clGetDeviceIDs和clGetDeviceInfo。用法和获取Platform信息类似。

cl_int clGetDeviceIDs(
	cl_platform_id platform, 
	cl_device_type device_type, 
	cl_uint num_entries, 
	cl_device_id* devices, 
	cl_uint* num_devices);

platform为所获取到的Platform ID
device_type为所要查询的device类型，有这么几种类型供选择：
CL_DEVICE_TYPE_CPU
CL_DEVICE_TYPE_GPU
CL_DEVICE_TYPE_ACCELERATOR
CL_DEVICE_TYPE_CUSTOM
CL_DEVICE_TYPE_DEFAULT
CL_DEVICE_TYPE_ALL
第一次调用该函数时，设置num_entries为0，devices设置为NULL，函数将通过num_devices返回所查询Platform中Device总数。
第二次调用该函数时，设置num_entries为刚才获取到的可用Device总数，函数通过devices返回Device ID的数组。

cl_int clGetDeviceInfo( 
	cl_device_id device, 
	cl_device_info param_name, 
	size_t param_value_size, 
	void* param_value,
	size_t* param_value_size_ret);

device设置为上一个函数获取到的Device ID数组中的项
param_name表示需要查询的具体信息，文档中给出了上百种类型，一般查询CL_DEVICE_NAME就可以了
param_value_size指明用于接受查询信息的数组长度
param_value则是接受查询信息的数组
param_value_size_ret是实际返回信息的长度

2.3 创建Context

    //step3: 创建Context
    cl_context_properties props[]={CL_CONTEXT_PLATFORM,(cl_context_properties)platform_ids[0],0};
    cl_context context;
    cl_uint context_param_value;
    size_t context_param_value_size_ret;
    
    context=clCreateContext(props,num_devices,device_ids,NULL,NULL,&ret);
    if(ret!=CL_SUCCESS) throw new OCLException(0,7,"clCreateContext Error");
    cout<<"Context: "<<context<<endl;

    ret = clGetContextInfo(context,CL_CONTEXT_REFERENCE_COUNT,4,&context_param_value,&context_param_value_size_ret);
    if(ret!=CL_SUCCESS) throw new OCLException(0,8,"clGetContextInfo Error");
    cout<<"Context Reference Count: "<<context_param_value<<endl;

上下文是用一个或多个设备创建的。OpenCL运行时使用上下文来管理命令队列、内存、程序和内核对象等，以及在上下文中指定的设备上执行内核。
其中，主要通过函数clCreateContext来创建Context，通过clGetContextInfo来查询Context信息。

cl_context clCreateContext(
	const cl_context_properties* properties,
	cl_uint num_devices,
	const cl_device_id* devices,
	void (CL_CALLBACK* pfn_notify)(const char* errinfo, const void* private_info, size_t cb, void* user_data), 
	void* user_data,
	cl_int* errcode_ret);

properties指定上下文属性名称及其相应值的数组。每个属性名（枚举类型）后面紧跟着相应的属性值，且数组以0结尾。支持的最重要的属性是CL_CONTEXT_PLATFORM，值设置为需要使用的Platform ID。
num_devices是在第一次调用clGetDeviceIDs时所获得的设备数。
devices是一个指针，指向第二次调用clGetDeviceIDs时获得的设备列表。
pfn_notify是一个可以由应用程序注册的回调函数。OpenCL将使用这个回调函数来报告在这个上下文中发生的错误信息。这个回调函数可以由OpenCL实现异步调用。应用程序负责确保回调函数是线程安全的。此回调函数的参数为：
errinfo：指向错误字符串的指针。
private_info和cb：指向OpenCL实现返回的二进制数据的指针，该数据可用于记录有助于调试错误的附加信息。
user_data：指向用户提供的数据的指针。
如果pfn_notify为空，则不注册回调函数。
user_data将作为user_data参数传递给pfn_notify函数。当pfn_notify为NULL时，user_data也为NULL。
errcode_ret将返回适当的错误代码。clCreateContext返回一个有效的context，如果context创建成功，errcode_ret设置为CL_SUCCESS。否则，它将返回一个空值，并在errcode_ret中返回以下错误值：
CL_INVALID_PLATFORM：如果属性为空且无法选择平台，或者在属性中指定的平台值不是有效的平台。
CL_INVALID_VALUE：如果属性中的上下文属性名不是受支持的属性名，或者为支持的属性名指定的值无效，或者多次指定同一属性名。
CL_INVALID_VALUE ：如果设备为空。
CL_INVALID_VALUE：如果num_devices等于零。
CL_INVALID_VALUE：如果pfn_notify为NULL，但user_data不为NULL。
CL_INVALID_DEVICE ：如果设备包含无效设备或未与指定平台关联。
CL_DEVICE_NOT_AVAILABLE：如果设备中的某个设备当前不可用，即使该设备是由clGetDeviceIDs返回的。
CL_OUT_OF_HOST_MEMORY ：如果在主机上分配OpenCL实现所需的资源失败。

cl_int clGetContextInfo(
	cl_context context, 
	cl_context_info param_name, 
	size_t param_value_size, 
	void* param_value,
	size_t* param_value_size_ret);

该函数与查询Platform和Device的函数类似，其中param_name的可查参数类型包括如下：

Context Info	Return Type	Description
CL_CONTEXT_REFERENCE_COUNT	cl_uint	Return the context reference count.
CL_CONTEXT_NUM_DEVICES	cl_uint	Return the number of devices in context.
CL_CONTEXT_DEVICES	cl_device_id[]	Return the list of devices and sub-devices in context.
CL_CONTEXT_PROPERTIES	cl_context_ properties[]	Return the properties argument specified in clCreateContext or clCreateContextFromType.If the properties argument specified in clCreateContext or clCreateContextFromType used to create context was not NULL, the implementation must return the values specified in the properties argument in the same order and without including additional properties.If the properties argument specified in clCreateContext or clCreateContextFromType used to create context was NULL, the implementation must return param_value_size_ret equal to 0, indicating that there are no properties to be returned.

2.4 创建Command Queue

OpenCL对象（如内存、程序和内核对象）是使用上下文创建的。对这些对象的操作是使用命令队列执行的。命令队列可用于按顺序对一组操作（称为命令）进行排队。跨多个命令队列共享对象需要应用程序执行适当的同步。

	//step4: 创建Command Queue
	cl_command_queue command_queue=clCreateCommandQueue(context, device_ids[0], 0, &ret);
	if(ret!=CL_SUCCESS) throw new OCLException(0,9,"clCreateCommandQueue Error");

其中通过函数clCreateCommandQueue创建命令队列。

cl_command_queue clCreateCommandQueue(
					cl_context context, 
					cl_device_id device, 
					cl_command_queue_properties properties, 
					cl_int *errcode_ret)

contex是前面创建的上下文。
device是前面获取的设备。
properties指明所创建的命令队列的属性，主要有CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ ENABLE属性值，若设置为该值，则命令队列无序执行。否则将顺序执行命令。
errcode_ret返回适当的错误代码。
若函数执行成功，则返回有效的非零命令队列，否则将返回空值。

2.5 创建Buffer

Buffer 由上下文 context 创建，这样上下文管理的多个设备就会共享 Buffer 中的数据。

	//step5: 创建Buffer
    const int ARRAY_SIZE=1000;
    const int BUFFER_SIZE=ARRAY_SIZE*sizeof(int);
    cl_mem buffer_A, buffer_B, buffer_C;
    int host_A[ARRAY_SIZE], host_B[ARRAY_SIZE], host_C[ARRAY_SIZE];

    for(int i=0;i<ARRAY_SIZE;i++)
    {
        host_A[i]=i;
        host_B[i]=i;
    }

	buffer_A=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR,BUFFER_SIZE,host_A,&ret);
	if(ret!=CL_SUCCESS) throw new OCLException(0,10,"clCreateBuffer Error");
	buffer_B=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR,BUFFER_SIZE,host_B,&ret);
	if(ret!=CL_SUCCESS) throw new OCLException(0,12,"clCreateBuffer Error");
	buffer_C=clCreateBuffer(context,CL_MEM_WRITE_ONLY,BUFFER_SIZE,NULL,&ret);
	if(ret!=CL_SUCCESS) throw new OCLException(0,13,"clCreateBuffer Error");
	cout<<"Buffer Object: "<<buffer_A<<"    "<<buffer_B<<"  "<<buffer_C<<endl;

其中，通过函数clCreateBuffer创建缓存对象。

cl_mem clCreateBuffer(
	cl_context context, 
	cl_mem_flags flags, 
	size_t size,
	void* host_ptr, 
	cl_int* errcode_ret);

context 是用于创建缓存对象的有效OpenCL上下文。
flags是一个位字段，用于指定内存对象所分配的位置和使用信息。下表描述了标志的可能值。
在创建缓存对象时，要考虑一下缓存对象的用途，如果是用来做数据传输的，将主存数据传输给显存数据，那么一般要设置CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR属性，如果内存对象是用来保存结果数据的，那么内存属性只需要设置CL_MEM_READ_WRITE属性。结果数据则需要在主机端通过：clEnqueueReadBuffer函数读取。
理论上Device端分配的内存运行速度更快，但实际应用中往往数据拷贝消耗性能比较多，而用Host端内存能做到类似“零拷贝”的效果，所以建议将内存分配在Host端。另外，有些OpenCL设备的Host与Device之间Cache交换做得比较好，所以内存对象在Device端还是Host端对运行速度的影响可能很小。

cl_mem_flags	描述
CL_MEM_READ_WRITE	这个标志指定缓存对象将由内核读写。这是默认值。
CL_MEM_WRITE_ONLY	这个标志指定缓存对象将被内核写入但不被读取。
CL_MEM_READ_ONLY	此标志指定缓存对象在内核中使用时是只读缓存对象。
CL_MEM_USE_HOST_PTR	此标志仅在host_ptr不为空时有效。如果指定，则表示应用程序希望OpenCL实现使用host_ptr引用的内存作为缓存对象的存储位。
CL_MEM_ALLOC_HOST_PTR	此标志指定应用程序希望OpenCL实现为缓存对象从主机可访问的内存（e.g. PCIe）分配存储位。CL_MEM_ALLOC_HOST_PTR和CL_MEM_USE_HOST_PTR是互斥的。
CL_MEM_COPY_HOST_PTR	此标志仅在host_ptr不为空时有效。如果指定，则表示应用程序希望OpenCL实现为缓存对象分配内存，并从host_ptr引用的内存中复制数据。CL_MEM_COPY_HOST_PTR和CL_MEM_USE_HOST_PTR是互斥的。CL_MEM_COPY_HOST_PTR可与CL_MEM_ALLOC_HOST_PTR一起使用，以初始化使用主机可访问的内存（e.g. PCIe）分配的cl_mem对象的内容。

size是要分配的缓冲区内存对象的大小（以字节为单位）。
host_ptr指向已由应用程序分配的缓冲区数据的指针。host_ptr指向的缓冲区大小必须大于等于size 字节。
errcode_ret将返回适当的错误代码。如果成功创建缓冲区对象，则errcode_ret设置为CL_SUCCESS。否则，它将返回一个NULL，并在errcode_ret中返回以下错误值之一：
CL_INVALID_CONTEXT：如果context 不是有效的上下文。
CL_INVALID_VALUE：如果flags 中指定的值无效
CL_INVALID_BUFFER_SIZE：如果size 为0或大于表4.3中为context 中所有设备指定的CL_DEVICE_MAX_MEM_ALLOC_size值。
CL_INVALID_HOST_PTR：如果host_ptr为空且CL_MEM_USE_HOST_PTR或CL_MEM_COPY_HOST_PTR：设置在标志中，或者如果host_ptr不为NULL，但在flags 中未设置CL_MEM_COPY_HOST_PTR或CL_MEM_USE_HOST_PTR。
CL_MEM_OBJECT_ALLOCATION_FAILURE：如果为缓冲区对象分配内存失败。
CL_OUT_OF_HOST_MEMORY：如果在主机上分配OpenCL实现所需的资源失败。

2.6 创建和编译OpenCL program

OpenCL program是使用OpenCL C语言编写的代码，其中实现了在GPU上执行的代码逻辑，需要通过clCreateProgramWithSource将代码文本加载到program对象中。Program对象与context相关联。Program对象导入代码后，还需要使用clBuildProgram对代码进行编译。到clBuildProgram这一步之前，所有的程序都是在CPU执行，这一步将OpenCL program构造成一个完整的、可加载到GPU上执行的程序。

	//step6: 创建OpenCL程序
	char fileName[]="./resource/vecadd.cl";
	char* kernel_code;
	int max_source_size=10000, source_size;
	FILE* fp=fopen(fileName,"r");
	if(!fp) throw new OCLException(0,14,"Failed to load kernel");
	kernel_code=(char*) malloc(max_source_size);
	source_size=fread(kernel_code,1,max_source_size,fp);
	fclose(fp);
	cl_uint kernel_code_count=1;
	char* kernel_codes[]={kernel_code};
	size_t lengths[]={source_size};
	
	cl_program program=clCreateProgramWithSource(context, kernel_code_count, kernel_codes, lengths, &ret);
	if(ret!=CL_SUCCESS) throw new OCLException(0,15,"clCreateProgramWithSource Error");
	ret=clBuildProgram(program, num_devices, device_ids, NULL, NULL, NULL);
	if(ret!=CL_SUCCESS) throw new OCLException(0,16,"clBuildProgram Error");

其中，clCreateProgramWithSource函数用于加载代码，clBuildProgram用于编译代码。

cl_program 	clCreateProgramWithSource (
					cl_context context, 	//关联的上下文
					cl_uint count,			//要加载的程序源码的个数
					const char **strings,	//程序源码指针列表，每个指针指向一段源码
					const size_t *lengths,	//每个程序源码对应的长度
					cl_int *errcode_ret) 	//错误码

cl_int clBuildProgram(cl_program program,	//上面创建的program
					  cl_uint num_devices, 	//设备数量
					  const cl_device_id *device_list,	//设备列表
					  const char *options, 	//编译选项，如头文件目录、库目录之类，没有可设NULL
					  void (*pfn_notify)(cl_program, void *user_data),	//异步执行时的回调函数
					  void *user_data)		//当回调函数不为空时，该参数可用于给回调函数传参数

构建program的OpenCL程序源码一般放在独立的文件中，后缀为.cl，通过读文件函数将文件内容读入字符串中。如下是本案例中用于计算向量相加的OpenCL程序源码。

#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable

__kernel void vecadd(__global int* A, __global int* B, __global int* C)
{
   int index=get_global_id(0);
   C[index]=A[index]+B[index];
}

__kernel限定符用来限定内核函数，且返回值必须为void。
__global限定变量在GPU的全局内存上，其它还有__constant常量全局内存，__local本地内存，在同一个work-group中可见，以及私有内存__private，所有未标明限定的都是私有变量，只在同一个work-item中可见。
get_global_id函数返回所请求维度上，执行该函数的work-item的全局索引。函数的参数指定第几维。
类似的内核函数还有：
get_work_dim()，返回NDRange的维度（1，2，3）
get_global_size(uint dimension)，返回所请求维度上work-item的总数
get_global_id(uint dimension)，返回所请求维度上执行该函数的work-item的全局索引。
get_local_size(uint dimension)，返回在所请求维度上work-group的大小
get_local_id(uint dimension)，返回在所请求维度上当前work-item在work-group中的索引
get_number_groups(uint dimension)，返回所请求维度上work-group的数目=global_size/local_size

2.7 创建kernel

加载到GPU上的program编译后，利用编译成功的program创建内核对象，并为内核对象的参数赋值。

	//step7: 创建kernel
	cl_kernel kernel=clCreateKernel(program, "vecadd", &ret);
	if(ret!=CL_SUCCESS) throw new OCLException(0,17,"clCreateKernel Error");
	
	ret=clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_A);
	if(ret!=CL_SUCCESS) throw new OCLException(0,18,"clSetKernelArg Error");
	ret=clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer_B);
	if(ret!=CL_SUCCESS) throw new OCLException(0,19,"clSetKernelArg Error");
	ret=clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_C);
	if(ret!=CL_SUCCESS) throw new OCLException(0,20,"clSetKernelArg Error");

其中，通过clCreateKernel函数创建内核对象，通过clSetKernelArg函数为内核对象的参数赋值。

cl_kernel clCreateKernel(cl_program   program,		//编译成功的OpenCL程序
						 const char   *kernel_name,	//是OpenCL程序中用__kernel限定符声明的函数名
		     			 cl_int       *errcode_ret) //错误码

cl_int clSetKernelArg(cl_kernel      kernel,	//内核对象
      				  cl_uint        arg_index,	//内核函数中的参数序号
            		  size_t         arg_size, 	
     				  const void     *arg_value)

arg_size指定参数的size。若参数类型为内存对象（memory object），则arg_size设定为sizeof(cl_mem)。若参数被声明为__local，则arg_size设定为需要为该本地对象分配的buffer的字节数。

2.8 执行kernel

	//step8: 执行kernel
	size_t dim = 1;
	size_t global_work_offset[] = {0};
	size_t global_work_size[] = {ARRAY_SIZE};
	size_t local_work_size[] = {64};
	ret=clEnqueueNDRangeKernel(command_queue, kernel, dim, NULL, global_work_size, NULL, 0, NULL, NULL);
	if(ret!=CL_SUCCESS) throw new OCLException(0,21,"clEnqueueNDRangeKernel Error");

OpenCL编程中的一个核心函数是clEnqueueNDRangeKernel，对于此函数的理解，有利于对数据在host和device之间的传递进行控制。

cl_int clEnqueueNDRangeKernel(cl_command_queue command_queue,	//有效的命令队列
							  cl_kernel kernel,					//有效的内核对象
							  cl_uint  work_dim,				//用于指定全局工作项和工作组中工作项的维度数
							  const size_t *global_work_offset, //当前版本必须为NULL，未来可用于描述工作项的全局偏移
							  const size_t *global_work_size,
							  const size_t *local_work_size, 
							  cl_uint num_events_in_wait_list,
							  const cl_event *event_wait_list,
							  cl_event *event)

global_work_size指向一个work_dim维度的无符号整数数组，这些值描述将执行内核函数的全局各维度中的work-item的数量。work-item的总数按global_work_size数组中各维的work-item数量的乘积计算。该总数不能超过设备的内存寻址大小，比如设备使用32位寻址空间，则global_work_size数组中各维度数值的乘积不能超过2^32.
local_work_size指向一个work_dim维度的无符号整数数组，该数组描述组成work-group的work-item数量。work-group中工作项的总数按local_work_size数组中各维的work-item数量的乘积计算。工作组中工作项的总数不能大于CL_DEVICE_MAX_WORK_GROUP_SIZE值（该值可通过终端下的clinfo命令查看，或通过clGetDeviceInfo函数查询），并且在local_work_size各维度中指定的work-item数必须小于或等于指定的CL_DEVICE_MAX_WORK_ITEM_SIZES各维度中的值。local_work_size 大小也可以是空值，在这种情况下，OpenCL将自动决定如何将全局工作项分解为适当的工作组实例。如果指定了local_work_size，则在global_work_size中指定的各维度work-item数必须可以被local_work_size对应维度中指定的work-item数整除。
event_wait_list和num_events_in_wait_list指定执行此特定命令之前需要完成的事件，可以为空。如果event_wait_list为空，则此特定命令不会等待任何事件完成，且num_events_in_wait_list必须为0。如果event_wait_list不为空，则其所指向的event_wait_list必须有效，且num_events_in_wait_list必须大于0。事件等待列表中指定的事件充当同步点。与event_wait_list和command_queue中的事件关联的上下文必须相同。
event返回标识此特定内核执行实例的事件对象。事件对象是唯一的，可以在以后用于标识特定的内核执行实例。如果event 为NULL，则不会为此内核执行实例创建任何事件，因此应用程序将无法查询或排队等待此特定的内核执行实例。

2.9 读取结果

	//step9: 读取结果
	clEnqueueReadBuffer(command_queue, buffer_C, CL_TRUE, 0, BUFFER_SIZE, host_C, 0, NULL, NULL);

	for(int i=0;i<ARRAY_SIZE;i++)
	{
		if(i>0&&i%16==0) cout<<endl;
		cout<<host_C[i]<<"    ";
	}

cl_int  clEnqueueReadBuffer(cl_command_queue command_queue, //有效命令队列
                 			 cl_mem buffer,					//有效的缓冲区对象
                             cl_bool blocking_read,			//阻塞（同步执行）还是非阻塞（异步执行）
                             size_t offset,					//缓冲区对象中要读取的偏移量（以字节为单位）
                             size_t cb,						//读取数据的字节数
                             void *ptr,						//指向主机内存中缓冲区的指针，数据将被读入其中。
                             cl_uint num_events_in_wait_list,
                             const cl_event *event_wait_list,
                             cl_event *event)				//如果是非阻塞调用，则event返回一个事件对象，可用于查询命令执行状态。

2.10 释放资源

	//step10: 释放资源
	clReleaseKernel(kernel);
	clReleaseProgram(program);
	clReleaseCommandQueue(command_queue);
	clReleaseMemObject(buffer_A);
	clReleaseMemObject(buffer_B);
	clReleaseMemObject(buffer_C);
	clReleaseContext(context);
	
	free(kernel_code);
	return 0;