c++ - 什么在 CUDA : global memory write + __threadfence() or atomicExch() to global memory? 中更快

Question

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

在代码中:

array[threadIdx.x] = value;

或者

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

谢谢。

Answer 1

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

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