【cuda编程】CUDA中的atomic原子操作

所谓原子操作，就是该操作绝不会在执行完毕前被任何其他任务或事件打断，也就说，它的 最小的执行单位，不可能有比它更小的执行单位。

CUDA的原子操作可以理解为对一个变量进行“读取-修改-写入”这三个操作的一个最小单位的执行过程。在它执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。基于这个机制，原子操作实现了在多个线程间共享的变量的互斥保护，确保任何一次对变量的操作的结果的正确性。

原子操作确保了在多个并行线程间共享的内存的读写保护，每次只能有一个线程对该变量进行读写操作，一个线程对该变量操作的时候，其他线程如果也要操作该变量，只能等待前一线程执行完成。原子操作确保了安全，代价是牺牲了性能。

1. 加法操作——atomicAdd()

读取位于全局或共享存储器中地址address 处的32 位或64 位字old，计算(old + val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。只有全局存储器支持64 位字。

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,unsigned long long int val);

代码举例：

#include <stdio.h>    
#include <stdlib.h>   
#include <cuda_runtime.h>  

__global__ void histo_kernel(unsigned int *histo)
{
    int atomic_value = atomicAdd(histo, 1);
    printf("atomic_value:%d, histo: %d\n", atomic_value, *histo);
}

int main(void)
{
	int threadSum = 3;

	//分配内存并拷贝初始数据
	unsigned int *dev_histo;

	cudaMalloc((void**)&dev_histo, sizeof(int));
	cudaMemcpy(dev_histo, &threadSum, sizeof(int), cudaMemcpyHostToDevice);

	histo_kernel <<<1,1 >>> (dev_histo);

	//数据拷贝回CPU内存
	cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);
	cudaFree(dev_histo);
	return 0;
}

输出：

atomic_value:3, histo: 4

2. 减法操作——atomicSub()

读取位于全局或共享存储器中地址address 处的32 位字old，计算(old - val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address, unsigned int val);

3. 交换操作——atomicExch()

读取位于全局或共享存储器中地址address 处的32 位或64 位字old，并将val 存储在存储器的同一地址中。这两项操作在一次原子事务中执行。该函数将返回old。只有全局存储器支持64 位字。

int atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address,unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address,unsigned long long int val);
float atomicExch(float* address, float val);

4. 最小值操作——atomicMin()

读取位于全局或共享存储器中地址address 处的32 位字old，计算old 和val 的最小值，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address,unsigned int val);

5. 最大值操作——atomicMax()

读取位于全局或共享存储器中地址address 处的32 位字old，计算old 和val 的最大值，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address,unsigned int val);

6. 增量操作——atomicInc()

读取位于全局或共享存储器中地址address 处的32 位字old，计算 ((old >= val) ? 0 : (old+1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

unsigned int atomicInc(unsigned int* address,unsigned int val);

7. 减量操作——atomicDec()

读取位于全局或共享存储器中地址address 处的32 位字old，计算 (((old == 0) | (old > val)) ? val : (old-1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

unsigned int atomicDec(unsigned int* address,unsigned int val);

8. 比较并交换——atomicCAS()

读取位于全局或共享存储器中地址address 处的32 位或64 位字old，计算 (old == compare ? val : old)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old（比较并交换）。只有全局存储器支持64 位字。

int atomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address,unsigned int compare,unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address,unsigned long long int compare,unsigned long long int val);

9. 与操作——atomicAnd()

读取位于全局或共享存储器中地址address 处的32 位字old，计算 (old & val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,unsigned int val);

10. 或操作——atomicOr()

读取位于全局或共享存储器中地址address 处的32 位字old，计算 (old | val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,unsigned int val);

11. 异或操作——atomicXor()

读取位于全局或共享存储器中地址address 处的32 位字old，计算 (old ^ val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,unsigned int val);