第九章 事件

本章中我们解释OpenCL事件以及如何使用事件。主要讨论一下内容：

• OpenCL中的基本事件模型
• 处理事件的API
• 用户定义的事件
• 用事件评测命令
• 内核中的事件

9.1 命令、队列和时间概述

命令队列是OpenCL的核心。平台定义了一个上下文，其中包含一个或多个计算设备。每个计算设备可以有一个或多个命令队列。提交到这些队列的命令完成OpenCL程序的具体工作。

在一个简单的OpenCL程序中，提交到一个命令队列的命令会按顺序执行。一个命令完成后，下一个命令才能开始，程序会展开为一个有严格顺序的命令序列。各个命令存在大量并发性，这种有序方法能够提供应用程序所需要的性能。

有一种常见的情况，即命令之间的依赖性能够表述为一个有向无环图。这种图可能包含独立的分支，可以安全、并发地运行。

应用程序有很多机会实现并发性，仅由命令队列是无法满足的。通过放宽这些顺序约束，在性能方面可能很有好处。不过，这些好处是有代价的。如果未使用命令队列的顺序语义来确保安全的命令执行顺序这个工作就要由程序员负责。在OpenCL中，这个任务可以利用事件来完成。

事件(event)是OpenCL中传递命令状态的对象。命令队列中的命令会生成事件，其他命令在执行之前可能要等待这些事件。

9.2 事件和命令队列

OpenCL事件是OpenCL中传递命令有关信息的对象。事件的状态描述了相关命令的状态。可以取以下状态值：

• CL_QUEUED:命令已经在命令队列中排队
• CL_SUBMITTED:入队的命令由宿主机提交给与命令队列关联的设备
• CL_RUNNING:计算设备正在执行命令
• CL_COMPLETE:命令已经完成
• ERROR_CODE:负值指示遇到某种错误条件。

创建事件有很多方法。最常见的时间源就是命令本身。在命令队列中排队的任何命令都会生成或等待事件。不同命令在API中都以同样的方式出现，所以我们可以利用一个例子来解释事件如何工作。考虑一个命令将内核入队，准备在一个计算设备上执行：

cl_int clEnqueueNDRangeKernel(cl_command_queue command_queue,
                              cl_kernel kernel,
                              cl_uint work_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  这个命令在执行之前需要等待完成的事件数
/*event_wait_list 这是一个指针数组，定义了这个命令等待的num_events_in_wait_list个事件。与event_wait_list中的事件和与command_queue关联的上下文必须相同。*/
/*cl_event* event：这是一个指针，指向这个命令生成的一个事件对象。可以由后续的命令或宿主机用来延续这个命令的状态*/　　

参数num_events_in_wait_list和*event_wait_list提供了合法值时，只有当列表中的所有事件都有CL_COMPLETE状态或者有一个负值指示某个错误条件时，命令才会运行。

一个例子

cl_event k_events[2];

err = clEnqueueNDRangeKernel(commands, kernel1, 1, NULL, &global, &local, 0, NULL, &k_events[0]);
err = clEnqueueNDRangeKernel(commands, kernel2, 1, NULL, &global, &local, 0, NULL, &k_events[1]);
err = clEnqueueNDRangeKernel(commands, kernel3, 1, NULL, &global, &local, 2, k_events, NULL);　　

这里有3个内核排队等待执行。前两个clEnqueueNDRangeKernel命令kernel1和kernel2入队。这些命令的最后一个参数会生成事件，事件将放在数组k_events[]的相应元素中。第三个clEnqueueNDRangeKernel命令kernel3入队。如clEnqueueNDRangeKernel的第7个参数和第8个参数所示，kernel3要等待数组k_events[]中的两个事件完成后，这个内核才会运行。不过，需要说明的是，将kernel3入队的最后一个参数时NULL。这表示我们不希望再生成一个让后续命令访问的事件。

如果需要详细控制命令执行的顺序，事件至关重要。不过，不需要这种控制时，可以让命令忽略事件(这包括事件的使用和生成)，这会很方便。可以使用下面的过程告诉一个命令忽略事件：

1. 将命令等待的事件数(num_events_in_wait_list)设置为0
2. 将事件数组的指针(*event_wait_list)设置为NULL。需要说明的是，如果设置这个指针为NULL，num_events_in_wait_list必须为0.
3. 将指向所生成事件的指针(*event)设置为NULL。

这个过程可以确保不会等待任何事件，也不会生成任何事件，这当然意味着对于这个特定的内核执行实例，应用程序不可能查询事件或将事件等待入队。

将命令入队时，命令需要指示一个同步点，该店之前的所有命令必须先完成，之后的命令的命令才能开始。可以使用clBarrier()函数对队列中的命令指示这样一个同步点。

cl_int clEnqueueBarrier(cl_command_queue command_queue)

cl_int clEnqueueMarker(cl_command_queue command_queue, cl_event* event)

cl_int clEnqueueWaitForEvents(cl_command_queue command_queue, cl_uint num_events, const cl_event* event_list)　　

这三个命令clEnqueueBarrier、clEnqueueMarker和clEnqueueWaitForEvents会对队列中的命令和同步点事假约束，这会影响OpenCL内存的一致性。它们共同为OpenCL中的同步协议提供基础构建模块。

9.3 事件对象　

事件都是对象。我们定义了3个函数来管理对象：

• clGetEventInfo　
• clRetainEvent
• clReleaseEvent

9.4 宿主机上生成事件

目前为止，事件都是由队列中的命令生成，影响同一个上下文队列中的其它命令。还可以使用事件来协调事件队列中运行的命令与宿主机上执行的函数之间的交互。首先来考虑宿主机上如何生成事件。可以通过在宿主机上创建用户事件来完成：

cl_event clCreateUserEvent(cl_context context, cl_int* errcode_ret)
//其中指定事件的上下文　　

返回的对象是一个事件对象，值为CL_SUBMITTED。这与OpenCL命令生成的事件相同，唯一的区别是用户事件由宿主机生成和管理。

对于命令队列中生成的事件，事件的状态由命令队列控制。不过，对于用户事件，事件的状态必须通过在宿主机上调用的函数显示控制。　

cl_int clSetUserEventStatus(cl_event event, cl_int exrcution_status)

clSetUserEventStatus只能调用一次来改变一个用户事件的执行状态，可以为CL_COMPLETE，或者改为一个负值指示存在一个错误。负整数值会导致等待这个用户事件的所有入队的命令终止。

9.5 影响宿主机执行的时间　　

//to do