OpenCL 伪随机生成器竞争条件

Question

以this question为基础，我实现了一个具有全局状态的伪随机数生成器：

  __global uint global_random_state;

  void set_random_seed(uint seed){
    global_random_state = seed;
  }

  uint get_random_number(uint range){
    uint seed = global_random_state + get_global_id(0);
    uint t = seed ^ (seed << 11);
    uint result = seed ^ (seed >> 19) ^ (t ^ (t >> 8));
    global_random_state = result; /* race condition? */
    return result % range;
  }

由于这些函数将在多个线程中使用，因此在写入 global_random_state 时会出现竞争条件。

这实际上可能有助于系统变得更加不可预测，所以这似乎是一件好事，但我想知道这是否有任何可能不会立即浮出水面的后果。当内核为运行?

时，GPU 内部是否存在任何可能导致问题的副作用

Answer 1

理论上，您需要 atom_cmpxchg 来确保正确性（或找到等效的 GPGPU）。然而，一个严重的警告，让整个机器通过一个缓存行序列化将从根本上扼杀你的性能。同一地址上的原子必须排成一个队列等待。不同位置的原子可以并行化（更多细节在最后）。

通常，在 GPGPU 上利用随机变量的算法会保留自己的随机变量生成器副本。这使每个工作项都可以缓存并可能重用它们自己的随机数，而不会在每个新随机数上使总线充满内存流量。搜索“OpenCL Monte Carlo”“模拟”或“示例”以获取样本。 CUDA 也有一些很好的例子。

另一种选择是使用随机生成器，它允许向前跳过并让不同的工作项按不同的顺序向前移动。虽然这可能需要更多的计算，但代价是您不会对内存层次结构造成太大的压力。

关于原子的更多血淋淋的细节：(1) GPU 缓存原子被设计为期望连续数组和原子 ALU 是每个存储区，(2) 缓存行中的每个双字每次都将由相同的原子 ALU 处理，并且(3) 相邻的缓存行将散列到不同的库。所以，如果每个时钟你都在连续的数据缓存行上做原子操作，那么工作应该完美地展开（或统计上如此）。相反，如果让每个工作项原子地修改相同的 32b，则缓存系统无法将所有相同的原子 ALU 插槽应用于 16/32/64（无论您的系统使用什么）。它必须将操作分解为 16/32/64 个单独的原子操作，以迭代方式应用它（通过上面的 #2）。在您有 512 个 ALU 来处理原子的系统中，您将每个时钟使用其中的一个 ALU（同一个）。把工作分散开来，你可以使用所有 512/c。

OpenCL 伪随机生成器竞争条件

OpenCL Pseudo random generator race condition

random

opencl

race-condition