OpenCL 死锁可能性

Question

我正在使用全局原子在 OpenCL 中的工作组之间进行同步。

所以内核使用类似的代码

... global volatile uint* counter;

if(get_local_id(0) == 0) {
    while(*counter != expected_value);
}
barrier(0);

等到counter变成expected_value。

在另一个地方确实如此

if(get_local_id(0) == 0) atomic_inc(counter);

理论上，如果所有工作组同时运行，该算法应该始终有效。但是如果一个工作组在另一个工作组完全完成后才开始，那么内核可能会死锁。

在 CPU 和 GPU（NVidia CUDA 平台）上，它似乎总是可以工作，有大量的工作组（超过 8000 个）。

对于算法，这似乎是最有效的实现。 （它对 2D 缓冲区中的每一行进行前缀求和。）

OpenCL 和/或 NVidia 的 OpenCL 实施是否保证这始终有效？

Answer 1

OpenCL 和/或 NVidia 的 OpenCL 实施是否保证这 总是有效？

就 OpenCL 标准而言，这并不能保证（对于 CUDA 也是如此）。现在，在实践中，由于您的特定 OpenCL 实施，它可能会很好地工作，但请记住，标准不保证，因此请确保您了解您的实施执行模型以确保这是安全的，并且此类代码不一定可以跨其他符合要求的实现移植。

从理论上讲，如果所有工作组同时运行，该算法应该始终有效

OpenCL 声明工作组可以按任何顺序运行，不一定是并行的，也不一定是并发的。 CUDA 有类似的措辞，尽管 CUDA 9 确实支持一种网格同步形式。

OpenCL 规范，3.2.2 执行模型：内核实例的执行：

符合要求的实现可以选择序列化工作组，因此正确的算法不能假设工作组将并行执行。没有安全且可移植的方式在工作组的独立执行之间进行同步，因为一旦进入工作池，它们就可以按任何顺序执行。