c++ - CUDA 中什么更快：全局内存写入 + __threadfence() 或 atomicExch() 到全局内存？

Question

假设我们有很多线程将顺序访问全局内存，那么总体上哪个选项执行得更快？我很怀疑，因为 __threadfence() 考虑了所有共享和全局内存写入，但写入是合并的。另一方面， atomicExch() 仅考虑重要的内存地址，但我不知道写入是否合并。

在代码中：

array[threadIdx.x] = value;

或者

atomicExch(&array[threadIdx.x] , value);

谢谢。

Answer 1

在 Kepler GPU 上，我敢打赌，atomicExch因为 Kepler 上的原子速度非常快。在 Fermi 上，这可能是一次洗涤，但鉴于您没有碰撞，atomicExch仍然可以表现良好。

请做一个实验并报告结果。

Answer 2

这两个人做的事情非常不同。

atomicExch确保没有两个线程一次尝试修改给定的单元格。如果发生这种冲突，一个或多个线程可能会停止。如果您事先知道没有两个线程访问同一个单元格，那么就没有必要使用任何atomic...函数了。

__threadfence()延迟当前线程（并且只有当前线程！）以确保给定线程的任何后续写入确实在以后发生。因此，__threadfence()就其本身而言，没有任何后续代码并不是很有趣。

因此，我认为没有必要比较这两者的效率。也许如果你能展示一个更具体的用例，我可以联系起来......

请注意，这些实际上都没有为您提供线程实际执行顺序的任何保证。