CUDA atomic原子操作

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

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

CUDA支持多种原子操作，常用的如下:

1、 atomicAdd()

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);

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

2、 atomicSub()

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

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

3、 atomicExch()

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);

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

4、 atomicMin()

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

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

5、 atomicMax()

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

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

6、 atomicInc()

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

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

7、 atomicDec()

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

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

8、 atomicCAS()

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);

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

9、 atomicAnd()

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

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

10、 atomicOr()

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

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

11、 atomicXor()

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

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

举个例子，定义1024个线程，求这1024个线程的ID之和，每个线程都会访问总和变量sum，如果不加原子操作，执行结果是错误并且是不确定的。

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

#define SIZE 1024

__global__ void histo_kernel(int size, unsigned int *histo)
{
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	if (i < size)
	{
		//*histo+=i;
		atomicAdd(histo, i);
	}
}

int main(void)
{
	int threadSum = 0;

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

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

	// kernel launch - 2x the number of mps gave best timing  
	cudaDeviceProp  prop;
	cudaGetDeviceProperties(&prop, 0);

	int blocks = prop.multiProcessorCount;
	//确保线程数足够
	histo_kernel << <blocks * 2, (SIZE + 2 * blocks - 1) / blocks / 2 >> > (SIZE, dev_histo);

	//数据拷贝回CPU内存
	cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);
	printf("Threads SUM：%d
", threadSum);
	getchar();
	cudaFree(dev_histo);
	return 0;
}

使用原子操作正确的结果是523776，不使用原子操作的结果不确定，其中一次执行结果是711，显然是不对的。