OpenCL 入队屏障与入队标记 - 有啥区别？

Posted 2023-03-17

【中文标题】OpenCL 入队屏障与入队标记 - 有啥区别？

【英文标题】：OpenCL enqueue barrier vs enqueue marker - what's the difference?OpenCL 入队屏障与入队标记 - 有什么区别？

【发布时间】：2021-08-09 22:17:15

【问题描述】：

OpenCL 规范具有 API 功能：

clEnqueueMarkerWithWaitList() clEnqueueBarrierWithWaitList()

_{（我正在查看 v2.2，但它们也存在于 v3.0 中。）}

它们不仅具有相同的签名，而且具有非常相似的描述。这两者有什么区别？

【参考方案1】：

我的回答假设 OpenCL 命令队列是有序队列的常见情况。乱序队列不太常见，需要更细粒度的同步。

查看标记和障碍之间区别的一种简单方法是查看它们各自最典型的用例场景。从异步程序的角度来看：

标记通常由

生产者

使用。

消费者

通常使用屏障。

因此，您通常会在一个队列上的命令在另一个队列上的命令完成之前无法开始执行的情况下使用它们。

示例： 假设您有 2 个命令队列：queue1 和 queue2。以及两个相关的内核：分别是 kernel1 和 kernel2。 kernel1 写入 kernel2 读取的缓冲区 (cl_mem)。

一系列命令记录（在单个线程上）可能如下所示：

queue1.enqueue(kernel1) 
event1 = queue1.enqueue_marker() // event 1 will signal only after kernel 1 is done
queue2.enqueue_barrier(wait_list = event1) // kernel 2 will not start before event1 is signaled
queue2.enqueue(kernel2)

您在实际代码中很少看到标记和屏障的原因是每个clEnqueue*** API 调用都隐含地充当标记和屏障。如果您无权访问原始的 clEnqueue*** 调用，则通常使用障碍和标记。

例如：

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, // barrier
    cl_event *event) // marker

返回的事件标记入队命令的结束（以及所有先前的命令，在有序队列的情况下）。您可以将返回的事件提供给不同队列（作为等待列表）上的消费者命令 (clEnqueue***) 以用作屏障。

【讨论】：

1.有时您需要等到某些命令执行后才开始另一个命令，因此即使您可以使用 clEnqueueWhatever，屏障也是有意义的。 2.您描述的标记和障碍之间的区别不是纯粹的约定问题吗？即在您的第一个代码 sn-p 中，我能否不简单地切换 enqueue_barrier 和 enqueue_marker 的使用并获得相同的行为？

1.为了等待一堆命令，您可以等待与最后一个命令关联的事件。或者将标记排入队列并等待返回的事件。 2. 您可能会使用标记的 waitlist 参数来获得类似屏障的行为。至少在一个有序队列上，同一队列上的命令之间存在隐式同步。没有等待名单的障碍只是为了标记一个事件似乎很奇怪。但我猜可能有用。