No.7_7 OpenCL 同步——命令队列

概述

内核代码的执行以及内存对象的操作，需要通过将 OpenCL 命令提交到命令队列来完成。在大多数情况下，我们只有一个命令队列，并且命令队列中的命令在执行时按照函数调用的顺序。但是在某些场景下并非如此，例如，当平台有多个 OpenCL 设备需要同时执行来提升性能；或者同一命令队列中的命令需要同时执行来提高并行度；抑或是应用程序需要跟踪命令的执行，来剖析程序的性能。这时候，为了保证命令的正确执行顺序，就需要在命令队列中、或者命令队列之间执行同步操作。

out-of-order VS in-order

调用 OpenCL 函数将命令提交到命令队列时，命令在命令队列中是按照函数调用顺序存放的。但是在命令执行时，其执行顺序并不一定和命令提交到命令队列的顺序一致。在创建命令队列的时候，可以通过设置 clCreateCommandQueue 函数的 properties 参数来指定命令队列中的命令是以 in-order 还是 out-of-order 的方式执行。

创建命令队列时，如果没有为命令队列设置 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 属性，提交到命令队列的命令将按照 in-order 的方式执行。在这种执行方式下，如果应用程序首先调用 clEnqueueNDRangeKernel 来执行内核 A，接着又调用 clEnqueueNDRangeKernel 来执行内核 B，可以认为内核 A 在 内核 B 之前完成。这时，如果内核 A 输出的内存对象作为内核 B 的输入，内核 B 将正确访问到内核 A 生成的数据。反之，如果在创建命令队列时设置了 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 属性，则不能保证内核 A 在内核 B 开始之前已经执行完成。

在创建命令队列时，为命令队列设置 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 属性，应用程序调用 OpenCL 函数把命令提交到命令队列后，将按照 out-of-order 的方式执行。在 out-of-order 执行模式下，不能保证提交到命令队列中的命令是按照函数调用的顺序执行。例如，上面通过调用 clEnqueueNDRangeKernel 来执行内核 A 和内核 B 的阐述中，内核 B 可能在内核 A 之前执行完成。为了保证两个内核以一个特定的顺序执行，可以使用事件同步机制。执行内核 A 时使用 event 事件对象来标识该命令，在随后调用 clEnqueueNDRangeKernel 执行内核 B 时，将 event 事件对象作为其 event_wait_list 参数成员。

除了事件同步可以用于多个内核执行时的同步，clEnqueueWaitForEvents（等待事件） 和 clEnqueueBarrier（屏障） 命令也可以提交到命令队列用于同步。等待事件命令保证在它之后的命令执行之前，之前提交到命令队列的命令（使用事件列表来标识）已经执行完成；屏障命令保证在