OpenCL编程指南-9.1命令、队列、事件

概述

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

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

不过，实际的应用程序通常并没有这么简单。大多数情况下，应用程序不需要严格有序地执行命令。执行其他命令的同时，内存对象可以在设备和宿主机之间移动。处理不相关内存对象的命令可以并发执行。在一般的应用程序中，同时运行的命令会提供充分的并发性。这种并发性可以由运行时系统加以利用，来增加可实现的并行性，从而得到显著的性能提升。

还有一种常见的情况，即命令之间的依赖性能够表述为一个有向无环图（Directed AcyclicGraph，DAG)。这种图可能包含独立的分支，可以安全、并发地运行。如果要求这些命令以串行顺序运行，就会对系统带来不必要的约束。无序的命令队列允许系统充分利用这些命令之间的并发性，不过还有更多并发性可以利用。对于可能与不同计算设备关联的不同命令队列，通过在这些命令队列上运行DAG的独立分支，可以利用大量额外的并发性。

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

事件（event）是OpenCL中传递命令状态的对象。命令队列中的命令会生成事件，其他命令在执行之前可能要等待这些事件。用户可以创建定制事件，在宿主机和计算设备之间提供额外的一层控制。事件机制可以用来控制OpenCL和图形标准（如OpenGL)之间的交互。最后，在内核中，程序员利用事件可以允许数据的移动与这些数据的相关操作重叠进行。

事件和命令队列

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

CL_QUEUED: 命令已经在命令队列中排队。
CL_SUBMITTED: 入队的命令由宿主机提交给与命令队列关联的设备。
CL_RUNNING: 计算设备正在执行命令。
CL_COMPLETE: 命令已经完成。
ERROR_CODE: 负值指示遇到某种错误条件。具体的值为平台或生成该事件的运行时API返回的值。

创建事件有很多方法。最常见的事件源就是命令本身。在命令队列中排队的任何命令都会生成或等待事件。不同命令在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)

现在，我们只对这个函数的最后三个参数感兴趣

cl_uint num_events_in_wait_list: 这个命令在执行之前需要等待完成的事件数。

const cl_event * 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状态或者有一个负值指示某个错误条件时，命令才会运行。

事件用来定义一个序列点，两个命令会在这里进入程序的一个已知状态，因此可以作为OpenCL中的一个同步点。与OpenCL中的所有同步点类似，根据OpenCL内存模型执行多个内核的情况下，内存对象会进入一个明确的状态。内存对象与一个上下文关联，所以即使计算涉及一个上下文中的多个命令队列，也能保证一致的状态。

例如，考虑下面这个简单的例子：

cl_event k_events[2];

//enqueue two kernels exposing events
err = clEnqueueNDRangeKernel(commands, kernel1, 1, NULL, &global, &lobal, 0, NULL, &k_events[0]);

err = clEnqueueNDRangeKernel(commands, kernel2, 1, NULL, &global, &lobal, 0, NULL, &k_events[1]);

//enqueue the next kernel..which waits for two prior
//events before launching the kernel
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_SUCCESS；否则，返回以下某个错误条件。

CL_INVALID_COMMMAND_QUEUE: 命令队列不是一个合法的命令队列。
CL_OUT_OF_RESOURCES: 在设备上分配OpenCL实现所需要的资源时失败。
CL_OUT_OF_HOST_MEMORY: 在宿主机上分配OpenCL实现所需要的资源时失败。

clEnqueueBarrier命令定义了一个同步点。这对于理解命令之间的顺序约束很重要。不过更重要的是，在介绍的OpenCL内存模型中，内存对象的一致性是针对同步点定义的。具体来说，在一个同步点上，命令可见的内存对象更新必须先完成，这样后面的命令才能看到新的值。

要定义更一般的同步点，OpenCL使用了事件和标志。标志用以下命令设置。

cl_int clEnqueueMarker(
    cl_command_queue command_queue,
    cl_event *event)

cl_command_queue command_queue: 应用这个标志的命令队列
cl_event *event: 这个指针指向用来传递标志状态的事件对象

只有当所有命令都入队之后，标志命令才能完成。对于一个有序队列，clEnqueueMarker命令的效果就类似于一个栅栏。但与栅栏不同的是，标志命令会返回一个事件。宿主机或其他命令可以等待这个事件，来确保标志命令完成前所有命令都已经入队。如果函数成功执行，clEnqueueMarker会返回CL_SUCCESS；否则，返回以下错误之一。

CL_INVALID_COMMAND_QUEUE: command_queue不是一个合法的命令队列。
CL_INVALID_VALUE: 事件是一个NULL值。
CL_OUT_OF_RESOURCES: 在设备上分配OpenCL实现所需要的资源时失败。
CL_OUT_OF_HOST_MEMORY: 在宿主机上分配OpenCL实现所需要的资源时失败。

下面这个函数将一个事件等待入队，这会等待一个特定的事件或-组事件完成，之后才会执行将来入队的命令。

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

cl_command_queue command_queue: 应用这个事件的命令队列
cl_uint num_events_in_wait_list: 这个命令等待完成的事件数
const cl_event *event_wait_list: 这是一个指针数组,定义了这个命令等待的num_events_in_wait_list个事件

这些事件定义了同步点。这意味着clEnqueueWaitForEvents完成时，内存模型中定义的内存对象更新必须完成，而且后续命令可以依赖于内存对象的一致状态。与event_list中的事件和与command_queue关联的上下文必须相同。

如果函数成功执行，则clEnqueuewaitForEvents返回CL_SUCCESS；否则，返回以下错误之一。

CL_INVALID_COMMMAND_QUEUE: command_queue不是一个合法的命令队列。
CL_INVALID_CONTEXT: 与command_queue和与event_list 中的事件关联的上下文不相同。
CL_INVALID_VALUE: num_events为0或event_list为NULL。
CL_INVALID_EVENT: event_list中指定的事件对象不是合法的事件。CL_OUT_OF_RESOURCES: 在设备上分配OpenCL实现所需要的资源时失败。CL_OUT_OF_HOST_MEMORY: 分配命令所需要的资源时失败。

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

例如，考虑两个队列共享同一个上下文，但是将命令指向不同的计算设备。可以在这两个设备之间共享内存对象（因为它们共享同一个上下文)，不过由于采用OpenCL的放宽的一致性内存模型，在某个给定点，共享的内存对象可能相对于一个队列（或另一个队列）中的命令处于一种不明确的状态。在一个策略点放置一个栅栏可以解决这个问题，程序员可能会用clEnqueueBarrier()命令来解决，如图9-1所示。
在这里插入图片描述
不过，OpenCL中的栅栏命令只会对该栅栏所在的命令队列施加约束，即约束其命令的顺序。程序员如何定义能够跨两个命令队列的栅栏呢? 如图9-2所示。

在其中一个队列中，clEnqueueMarker()命令入队，返回一个合法的事件对象。标志就相当于它自己的队列中的一个栅栏，不过还会返回一个事件，其他命令可以等待这个事件。在第二个队列中，我们在期望的位置设置了一个栅栏，并在栅栏后面增加了一个clEnqueueWaitForEvents调用。clEnqueueBarrier命令会使相应队列有期望的行为，也就是说，clEnqueueBarrier()之前的所有命令必须先完成，后面的命令才能执行。clEnqueueWaitForEvents()调用可以定义从其他队列到标志的连接。最终结果是得到一个同步协议，可以在两个队列之间定义栅栏功能。