带有 clSetKernelArg 的 OpenCL 竞争条件

Posted 2023-03-29

【中文标题】带有 clSetKernelArg 的 OpenCL 竞争条件

【英文标题】：OpenCL race condition with clSetKernelArg

【发布时间】：2020-11-22 12:27:56

【问题描述】：

来自 Khronos 网站上关于 clSetKernelArg 的线程安全：

所有 OpenCL API 调用都是线程安全的，除了 clSetKernelArg，它可以安全地从任何主机线程调用，并且只要并发调用对不同的 cl_kernel 对象进行操作，就可以安全地重新调用。但是，如果同时从同一个 cl_kernel 对象上的多个主机线程调用 clSetKernelArg，则 cl_kernel 对象的行为是不确定的。

我的问题是，有没有办法定义这种行为，内核可以从多个线程中读取和写入单个内核对象？

我认为被内核修改的对象上的 std::atomic 会阻止这种未定义的行为，但根据我的尝试，它会导致内核的输出产生错误的值。有没有更好的方法来实现这种/已知的处理案例的技术？

在分配对象的大小如此之大以至于为每次内核执行重新创建一个新对象会占用太多内存的情况下，这可能很有用，并且首选共享/可覆盖对象。

【问题讨论】：

您确定完全清楚 OpenCL 标准对您提出的要求吗？这似乎并不不合理或难以确保。请注意，它仅适用于 clSetKernelArg() 调用本身；您仍然可以在不同的线程上填充缓冲区（实际上，如果它是内存映射的，则在多个主机线程上填充单个缓冲区）。

这只是 clSetKernelArg 和 clEnqueueNDRangeKernel 之间的竞争条件，所以我只是使用每个内核的互斥锁来防止它。

【参考方案1】：

内核可以从多个线程中读取和写入单个内核对象？

“内核”是指在 GPU 上执行的代码片段，而“单个内核对象”是指主机代码中的 cl_kernel ？ GPU 上的内核永远不会看到存在于主机端的 cl_kernel 结构。我假设您说的是内核使用缓冲区对象 (cl_mem) 参数。

你可以把cl_kernel想象成：

struct 
  size_t num_args;
  void* args[];
 _cl_kernel;
typedef struct _cl_kernel * cl_kernel;

如果您调用 clSetKernelArg()，它只会在该结构中设置一些内容。如果您调用 clEnqueueNDRangeKernel()，它会获取 cl_kernel 结构（参数）的快照，并将其附加到某个内部设备队列。 “快照”并不是说它创建了一个实际 cl_mem 缓冲区内容的隐藏快照；它只是复制对 cl_mem 参数的引用。由于它是一个引用，因此无论您是使用来自多个线程的单个 cl_kernel 对象，还是使用相同名称多次调用 clCreateKernel，然后使用那些 cl_kernel 在每个线程中；这只是方便的问题，最终结果是一样的。

如果您有一个有序命令队列，您的内核将按入队顺序确定性地执行。如果您有多个命令队列（有序或无序，没关系），则没有任何隐式排序队列之间，所以如果你将相同的内核排入所有队列，它们将以随机顺序执行。您可以使用事件强制显式顺序。 IOW，你做到了：

cl_event event1, event2;
cl_kernel K;
...
clEnqueueNDRangeKernel(queue_1, K, ... , &event1);
clEnqueueNDRangeKernel(queue_2, K, ... , 1, &event1, &event2);

等等。这将强制内核执行等待前一个，即使它们在不同的队列中。但是一次只能有一个内核使用缓冲区。

如果您希望多个正在运行的内核同时使用同一个缓冲区，那么这取决于该缓冲区的使用模式。如果您只进行读取，则可以安全地同时使用来自任意数量内核的缓冲区。对于写入用途，如果您知道您只会写入缓冲区的一部分，您可以尝试使用子缓冲区 (clCreateSubBuffer)。否则你可能不走运（也许你可以尝试原子操作，但它可能会使算法变得非常慢）。