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) 缓存行中的每个 dword 每次都将由同一个原子 ALU 处理，以及 (3 ) 相邻的缓存线将散列到不同的银行。因此，如果每个时钟都在连续的数据缓存行上进行原子操作，那么工作应该完美地展开（或统计上如此）。相反，如果让每个工作项以原子方式修改相同的 32b，则缓存系统无法将所有相同的原子 ALU 插槽应用于 16/32/64（无论您的系统使用什么）。它必须在 16/32/64 个单独的原子操作中分解操作，并迭代地应用它（通过上面的 #2）。在一个有 512 个 ALU 来处理原子的系统中，每个时钟都将使用其中的 1 个（同一个）。展开工作，您可以使用所有 512/c。