CUDA编程之ATOM操作

[_梦魇花葬]

原子函数（atomic function）对位于全局或共享存储器的一个32位或64位字执行read-modify-write的原子操作。也就是说，当多个线程同时访问全局或共享存储器的同一位置时，保证每个线程能够实现对共享可写数据的互斥操作；在一个操作完成之前，其他任何线程都无法访问此地址。例如下面的一个程序：
__global__ void testAtom(int * g_odata)
{
const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

//算数原子操作
atomicAdd(&g_odata[0], 10); // Atomic addition
atomicMax(&g_odata[0], tid); //Atomic maximum
atomicInc((unsigned int *)& g_odata[0], 17); //Atomic increment
atomicCAS(&g_odata[0], tid-1, tid); //Atomic compare - and - swap
//位原子操作
atomicAnd(&g_odata[8], 2 * tid +7); // Atomic AND
atomicOr(&g_odata[8], 1<< tid); // Atomic OR
atomicXor(&g_odata[8], tid); //Atomic XOR
}
如果上述程序不使用原子操作，则同一half warp内的线程读入数据进行运算后，将同时试图改变存储器中的值，存储器中的结果可能是其中任意一个线程的结果。
备注：原子操作仅使用于有符号和无符号整型，但atomicExch除外，它还支持单精浮点。

[linxxx3]

补充一下，原子操作的主要场景，是多个SM上的线程同时对同一显存地址进行操作（至少一个是写操作）。产生这种情况，一般访问地址由输入数据计算得到，无法通过设计算法避免地址冲突；否则还是应该尽量避免使用原子操作。 一个warp内原子操作，也可以操作同一地址，不过非常影响执行速度。

[maowrx]

新人入门学习中

[_梦魇花葬]

很对的！原子操作很复杂的，用的时候要分清当时的情况，不要盲目的去使用，有的时候不适合用原子操作。还有不同的硬件，原子操作的类型也不一样，都需要一个系统的学习！大家互相学习！

[GW786228836]

[xzzy111]

感谢分享00