OpenCL 内核中的读写竞争条件会导致数据损坏吗？答案

【问题标题】：Can read-write race conditions in an OpenCL kernel lead to corrupted data?OpenCL 内核中的读写竞争条件会导致数据损坏吗？
【发布时间】：2015-05-21 20:37:07
【问题描述】：

考虑一对读取和写入相同内存位置的 OpenCL 内核。作为一个简单的例子，考虑以下 OpenCL 程序：

__kernel void k1(__global int * a)
{
    a[0] = 2*a[1];
}

__kernel void k2(__global int * a)
{
    a[1] = a[0]-1;
}

如果启动多个线程，运行其中的多个内核，则全局内存的最终状态是不确定的。

这仍然可能允许人们编写异步算法来接受内核中任何可能的操作顺序。

但是，这要求对全局 GPU 内存的读写是原子的。

我的问题是

这是否保证在任何当前的 GPGPU 硬件上都是正确的？
如果这被 OpenCL 标准视为未定义的行为？如果是这样，常见的实现（特别是包含在 CUDA 工具包中的）有什么作用？
如何测试这种担忧？

【问题讨论】：

为什么会有两个同名的内核？无论如何，内存读取不是原子的。
谢谢，错字已修正。你能更精确一点吗？ “内存读取不是原子的”到底是什么意思？ OpenCL 标准不能保证它们是原子的吗？它们在任何现有硬件上都不是原子的？你能指出某种参考吗？

标签： opencl gpu gpgpu nvidia

【解决方案1】：

如果您将内核命令排入单个命令队列，该队列被创建为 in-order 队列（即您在创建它时没有指定CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE），那么只有一个内核命令将一次执行。这意味着您不会在不同内核实例之间遇到任何此类问题（尽管您仍然可能在单个内核实例中的工作项之间存在竞争条件，如果它们正在访问相同的内存位置）。

如果您使用无序队列或多个命令队列，那么您可能确实存在竞争条件。无法保证您的加载-修改-存储序列将是原子操作，这将导致未定义的行为。

根据您实际想要对内核执行的操作，您可以使用 OpenCL 的内置原子函数，这些函数允许您以原子方式执行一组特定的读-修改-写操作.

【讨论】：

我的问题不是关于 load-modify-store 序列是原子的，而是读写本身是原子的。