OpenCL伪随机生成器竞争条件 [英] OpenCL Pseudo random generator race condition

问题描述

以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内部是否有可能在以后运行内核时导致问题的副作用？

gpgpu

理论上您希望atom_cmpxchg在这里正确无误(或者找到等价的推荐答案)。然而，一个严重的警告是，让整个机器通过单个缓存线串行化将从根本上扼杀您的性能。同一地址上的原子必须形成队列并等待。不同位置的原子可以并行(更多详细信息见末尾)。

通常，在GPGPU上利用随机变量的算法将保留其自己的随机变量生成器副本。这使得每个工作项都可以缓存并潜在地重用它们自己的随机，而不必在每个新的随机上使用内存流量使总线不堪重负。搜索"；OpenCL Monte Carlo&或"；Simulation"；或"；Example&Quot；以获取示例。库达也有一些很好的例子。

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

有关原子的更多血腥细节：(1)GPU缓存原子被设计为期望连续数组和原子ALU是每个存储体的，(2)缓存行中的每个双字每次都将由相同的原子ALU处理，以及(3)相邻的缓存行将散列到不同的存储体。因此，如果每个时钟都在连续的数据缓存线上进行原子化，那么工作应该是完全分散的(或者在统计上是这样)。相反，如果让每个工作项自动修改相同的32b，那么缓存系统就不能将所有相同的原子ALU槽应用到16/32/64(无论您的系统使用什么)。它必须在16/32/64中将操作拆分，单独的原子操作迭代地应用它(通过上面的#2)。在有512个ALU来处理原子的系统中，每个时钟(相同的时钟)将使用这些ALU中的1个。展开工作，您就可以使用所有512/c。

这篇关于OpenCL伪随机生成器竞争条件的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！