OpenCL 内核随机写入 __global 缓冲区是否安全？答案

【问题标题】：Is it safe for an OpenCL kernel to randomly write to a __global buffer?OpenCL 内核随机写入 __global 缓冲区是否安全？
【发布时间】：2019-11-08 20:20:00
【问题描述】：

我想运行一个经过检测的 OpenCL 内核来获取一些执行指标。更具体地说，我添加了一个隐藏的全局缓冲区，它将使用 N 个零从主机代码初始化。 N 个值中的每一个都是整数，它们代表不同的度量，每个内核实例将以不同的方式递增，具体取决于其执行路径。

一个简单的例子：

__kernel void test(__global int *a, __global int *hiddenCounter) {
    if (get_global_id(0) == 0) {
        // do stuff and then increment the appropriate counter (random numbers here)
        hiddenCounter[0] += 3;
    }
    else {
        // do stuff...
        hiddenCounter[1] += 5;
    }
}

内核执行完成后，我需要主机代码聚合（一个简单的元素向量加法）所有hiddenCounter 缓冲区并打印适当的结果。

我的问题是当多个内核实例尝试写入hiddenCounter 缓冲区的同一索引时是否存在竞争条件（这肯定会在我的项目中发生）。我需要强制执行某种同步吗？或者这对于__global 参数是不可能的，我需要将其更改为__private？之后我能否从主机代码中聚合 __private 缓冲区？

【问题讨论】：

标签： memory opencl

【解决方案1】：

我的问题是当多个内核实例尝试写入hiddenCounter 缓冲区的同一索引时是否存在竞争条件

答案是肯定是，您的代码将容易受到当前编写的竞争条件的影响。

我需要强制执行某种同步吗？

是的，您可以为此目的使用global atomics。除了最古老的 GPU 之外，所有的 GPU 都将支持这一点。（任何支持 OpenCL 1.2 或 cl_khr_global_int32_base_atomics 和类似扩展的东西）

请注意，这将产生不小的性能开销。根据您的访问模式和频率，在private 或local 内存中收集中间结果并将它们写入内核末尾的全局内存可能会更快。（在local 的情况下，整个工作组将只为每个更新的单元共享一个全局原子调用——不过，您需要使用局部原子或缩减算法来累积整个组中各个工作项的值。）

另一种选择是使用更大的全局内存缓冲区，为每个工作项或组配备计数器。在这种情况下，您不需要原子来写入它们，但您随后需要组合主机上的值。这显然会使用更多的内存，而且可能还会使用更多的内存带宽——现代 GPU 应该缓存对您的小 hiddenCounter 缓冲区的访问。所以你需要找出/尝试在你的情况下哪个是较小的邪恶。

【讨论】：