共享内存上的原子操作

Question

如何对共享内存进行原子操作？
我有类似的东西：

__shared__ int a[10];
//set a
if(tid<5)
  a[2]++;

因此 5 个线程正在递增 a。我该怎么做？
我知道以这种方式我正在序列化 5 个线程的执行，但这对扭曲有何影响？ warp 中的所有线程会被序列化还是只序列化前 5 个？

Answer 1

将 a[2]++ 替换为

atomicAdd(a + 2, 1);

如果你可以做一个无符号的，你可能更喜欢使用 atomicInc() 代替，但任何一个都会影响性能。