任务队列允许程序员通过编辑直接创建、管理和调度任务。队列中的命令按提交命令的顺序执行，队列中的吓一跳命令必须等待当前命令执行完后才能开始执行。不同命令队列的先后执行顺序得不到保证。
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
当不需要对命令精确控制时，可设置参数来忽略事件对象
- 将 num_events_in_wait_list 设为0
- 将 *event_wait_list 设为NULL
- 将 *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高级特征被砍掉啦~
