OpenCL并行编程基础 第六章 OpenCL事件和队列

任务队列允许程序员通过编辑直接创建、管理和调度任务。队列中的命令按提交命令的顺序执行，队列中的吓一跳命令必须等待当前命令执行完后才能开始执行。不同命令队列的先后执行顺序得不到保证。

6.1 命令、命令队列和事件

命令队列保证FIFO：

/*
    启动顺序任务队列的代码
*/
cl_uint num_devices;
cl_device_id devices[1];
errNum = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_CPU, 1, &devices[0], &num_devices);
context = clCreateContext(0, 1, devices, NULL, NULL, &errNum);
cl_command_queue queue_cpu;
queue_cpu = clCreateCommandQueue(context, devices[0], 0/* 顺序执行 */, &errNum);
/*向队列中插入任务*/

乱序队列：

queue_cpu = clCreateCommandQueue(context, devices[0], CL_QUEUE_OUT_OF_ORDER_EXEC_MEOD_ENABLE,&errNum);

单一队列会导致大量计算资源没有用到，故可以使用多队列机制。：

cl_uint num_devices;
cl_devices_id devices[2];

errNum = clGetDevicesIDs(NULL, CL_DEVICES_TYPE_CPU, 1, &devices[0], &num_devices);
errNum = clGetDevicesIDs(NULL, CL_DEVICES_TYPE_GPU, 1, &devices[1], &num_devices);
context_cpu = clCreateContext(0, 1, &devices[0], NULL, NULL, &errNum);
context_gpu = clCreateContext(0, 1, &devices[1], NULL, NULL, &errNum);
cl_command_queue queue_cpu,queue_gpu;
queue_cpu = clCreateCommandQueue(context_cpu, devices[0], 0, &errNum);
queue_gpu = clCreateCommandQueue(context_gpu, devices[1], 0, &errNum);
//CPU、GPU的两个队列对应两个自己的上下文，互相不可见，也不存在sync，常会出错

上面这个问题要将CPU、GPU放在一个上下文中，然后两个队列的任务统一派送到该上下文中，但会出现内存访问不能进行同步的问题。

cl_uint num_devices;
cl_devices_id devices[2];

errNum = clGetDevicesIDs(NULL, CL_DEVICES_TYPE_CPU, 1, &devices[0], &num_devices);
errNum = clGetDevicesIDs(NULL, CL_DEVICES_TYPE_GPU, 1, &devices[1], &num_devices);
context = clCreateContext(0, 2, devices, NULL, NULL, &errNum);
cl_command_queue queue_cpu,queue_gpu;
queue_cpu = clCreateCommandQueue(context, devices[0], 0, &errNum);
queue_gpu = clCreateCommandQueue(context, devices[1], 0, &errNum);
//CPU、GPU的两个队列对应两个自己的上下文，互相不可见，也不存在sync，常会出错

所以引入事件机制。

6.2 事件的定义与基本用法

事件（event）是命令相关的数据对象，表示该命令的状态，它的列表

命令状态

状态取值	含义
CL_QUEUED	该命令已经进入了队列
CL_SUBMITTED	该命令已经被主机提交到了该队列对应设备上
CL_RUNNING	该设备正在执行该命令
CL_COPMLETE	该命令已经执行完毕
ERROR_CODE	该命令执行过程中出现了错误

 创建方法，常见的事件来源命令本身

//API原型函数
cl_int clEnqueueNDRangeKernel(
    cl_command_queue commandqueue,
    cl_kernel kernel,
    cl_uint word_dim,
    const size_t *global_work_offset,
    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
)

事件相关参数

形参	含义
num_events_in_wait_list	该命令要等待的事件的数量
evevt_wait_list	该命令要等待时间的列表
event	该命令所产生的事件

当非零的 num_events_in_wait_list 和有效的 event_wait_list 值传给API时，相应命令将等待至 event_wait_list 中所有的 num_events_in_wait_list 个事件均为CL_COMPLITE或ERROR_CODE时才会执行。 实例：

cl_event events[2];
errNum = clEnqueueNDRangeKernel(commands, kernel_A, 1, NULL, &global, &local, 0, NULL, &events[0]);//无等待
errNum = clEnqueueNDRangeKernel(commands, kernel_B, 1, NULL, &global, &local, 0, NULL, &events[1]);//无等待
errNum = clEnqueueNDRangeKernel(commands, kernel_C, 1, NULL, &global, &local, 2,events, NULL);//等待kernel_A,kernel_B

当不需要对命令精确控制时，可设置参数来忽略事件对象

1. 将 num_events_in_wait_list 设为0
2. 将 *event_wait_list 设为NULL
3. 将 *event 设为NULL

同步点（执行这条命令前要保证提交的所有命令已执行完毕）查询方法

cl_int clEnqueueBarrier(cl_command_queue command_queueu)
//只接受同步点关联的队列，clEnqueueBarrier通过队列来阻塞函数

用户可以通过一个事件来表示队列中命令的完成状态

cl_int clEnqueueMaker(cl_command_queue commandqueue, cl_event *event)
//commandqueue是同步点要提交到的队列，event是同步点要生成的函数，在event变成CL_COMPLETE前，所有等待次事件的命令都将等待

还有用户显示地等待一个或多个事件 clEnqueWaitForEvents(...)

6.3 事件对象与用户事件

事件对象本身的性质，通过三个函数来管理，clGetEventInfo 查询事件对象存放的状态、clRetainEvent 查询事件对象的引用数、clReleaseEvent 减少一个事件对象的引用数;

如果用户事件创建成功，啧clCreateUserEvent函数将返回一个状态为CL_SUBMITTED的事件对象。

cl_event clCreateUserEvent(cl_context context,cl_int *errcode_ret)
cl_int clSetUserEventStatus(cl_event event, cl_int execution_status)
cl_int clWaitForEvents(cl_uint num_events, const cl_event *event_list)

6.5 使用事件进行性能剖析例子

几乎所有命令都可以在产生相应队列时设置CL_QUEUE_PROFILING_ENABLE标志，即可使用clGetEventProfilingInfo函数获得执行时间。

/* Profiling APIs */
extern CL_API_ENTRY cl_int CL_API_CALL
clGetEventProfilingInfo(cl_event            /* event 要剖析的事件 */,
                        cl_profiling_info   /* param_name 计算时间的方式,
                        size_t              /* param_value_size 剖析数据的预期大小*/,
                        void *              /* param_value 返回数据的指针*/,
                        size_t *            /* param_value_size_ret 剖析数据的实际大小尺寸*/) CL_API_SUFFIX__VERSION_1_0;

cl_profiling_info

取值	含义
CL_PROFILING_COMMAND_QUEUE	主机将命令放入队列的时间，单位为纳秒
CL_PROFILING_COMMAND_SUBMIT	命令提交的时间
CL_PROFILING_COMMAND_START	命令开始执行的时间
CL_PROFILING_COMMAND_END	命令结束执行的时间

思路：首先产生命令队列，设置为运行剖析的形式，然后放别获得命令进入队列和执行的时间，然后两个做差

cl_command_queue commandQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, NULL);//开启运行剖析
cl_event event;
errNum = clEnqueueNDRangeKernel(commandQueue, kernel,2, global_work_offset,globalWorkSize, localWorkSize,0, NULL, &event);
clFinish(commandQueue);
errNum = clWaitForEvents(1,&event);
cl_ulong start, end; double runTime;
size_t return_bytes;
errNum = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong),&start,&return_bytes);
errNum = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, &return_bytes);
runTime = (double)(end - start)*1.0e-9;
std::cout << "commandQueue从开始到结束运行时间：" <<runTime<< " 秒"<<std::endl;

PS：我用的是GTX 1060，不支持OpenCL2.0 chapter7 OpenCL2.0高级特征被砍掉啦~