OpenCL内核随机写入__global缓冲区安全吗？
我想运行一个工具化的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缓冲区？
回答 1
Stack Overflow用户
发布于 2019-11-08 12:20:00
我的问题是，当多个内核实例试图写入
hiddenCounter缓冲区的同一个索引时，是否存在争用条件？
这个问题的答案是强调是，您的代码将很容易受到当前编写的种族条件的影响。
我需要强制执行某种同步吗？
是的，您可以为此目的使用全局原子。除了最古老的GPU之外，所有的GPU都会支持这一点。(任何支持OpenCL 1.2或cl_khr_global_int32_base_atomics及类似扩展的内容)
请注意，这将有一个重要的性能开销。根据您的访问模式和频率，收集private或local内存中的中间结果并将它们写入内核末尾的全局内存可能会更快。(在local的情况下，整个工作组将只为每个更新的单元共享一个全局原子调用-您将需要使用本地atomics或简化算法来在整个组中累积来自单个工作项的值。)
另一种选择是使用更大的全局内存缓冲区，为每个工作项或组设置计数器。在这种情况下，您将不需要atomics对它们进行写入，但随后需要将主机上的值组合起来。显然，这会使用更多的内存，而且可能也会占用更多的内存带宽--现代GPU应该缓存对小型hiddenCounter缓冲区的访问。因此，你需要解决/尝试哪一个在你的情况下是较小的邪恶。
https://stackoverflow.com/questions/58757657
复制相似问题
腾讯云开发者
