在CUDA中，哪个更快：全局内存写入+__threadfence()还是对全局内存的atomicExch()操作？

Question

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

代码如下：

array[threadIdx.x] = value;

或者

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

谢谢。

- dsilva.vinicius

---

可以尝试两种方法并报告结果？ :) 或者，可以尝试两种方法，然后惊喜地问：“但为什么ABC比XYZ要快那么多？”或者，尝试两种方法，看到平淡无奇的结果然后继续前进。 - user166390

Answer 1

在Kepler GPU上，我会选择使用 atomicExch ，因为原子操作在Kepler上非常快。在Fermi上，可能会相差无几，但是考虑到没有冲突，atomicExch 仍然可以表现出良好的性能。

请进行一次实验并报告结果。

- harrism

Answer 2

这两个函数的作用非常不同。

atomicExch 确保没有两个线程同时尝试修改给定单元格。如果出现这样的冲突，一个或多个线程可能会被阻塞。如果您事先知道没有两个线程访问同一单元格，则没有必要使用任何atomic...函数。

__threadfence() 延迟当前线程（仅当前线程！）以确保随后由给定线程编写的任何写入实际上稍后发生。 因此，__threadfence()本身，没有任何后续代码，是不太有趣的。

出于这个原因，我认为没有比较这两者的效率的必要性。也许如果您能展示更具体的用例，我就可以相关联……

请注意，这两个函数都不能保证线程的实际执行顺序。

- CygnusX1