CUDA中更快的是:全局内存写入+__threadfence(）或atomicExch(）到全局内存

What is faster in CUDA: global memory write + __threadfence() or atomicExch() to global memory?

本文关键字：全局 内存 atomicExch threadfence CUDA      更新时间：2023-10-16

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

代码中：

array[threadIdx.x] = value;

或

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

谢谢。

在开普勒GPU上，我会打赌atomicExch，因为原子学在开普勒上非常快。在费米上，这可能是一次洗涤，但如果没有碰撞，atomicExch仍然可以表现得很好。

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

这两者做的事情截然不同。

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

__threadfence()延迟当前线程（并且仅延迟当前线程！），以确保给定线程的任何后续写入都会在以后发生。因此，__threadfence()本身在没有任何后续代码的情况下并不是很有趣。

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

请注意，这两种方法实际上都不能保证线程的实际执行顺序。

相关文章：