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，如果不加原子操作，执行结果是错误并且是不确定的。

[cpp]  view plain  copy

 print ?

1. #include <stdio.h>      
2. #include <stdlib.h>     
3. #include <cuda_runtime.h>    
4.   
5. #define SIZE 1024  
6.   
7. __global__ void histo_kernel(int size, unsigned int *histo)  
8. {  
9.     int i = threadIdx.x + blockIdx.x * blockDim.x;  
10.     if (i < size)  
11.     {  
12.         //*histo+=i;  
13.         atomicAdd(histo, i);  
14.     }  
15. }  
16.   
17. int main(void)  
18. {  
19.     int threadSum = 0;  
20.   
21.     //分配内存并拷贝初始数据  
22.     unsigned int *dev_histo;  
23.   
24.     cudaMalloc((void**)&dev_histo, sizeof(int));  
25.     cudaMemcpy(dev_histo, &threadSum, sizeof(int), cudaMemcpyHostToDevice);  
26.   
27.     // kernel launch - 2x the number of mps gave best timing    
28.     cudaDeviceProp  prop;  
29.     cudaGetDeviceProperties(&prop, 0);  
30.   
31.     int blocks = prop.multiProcessorCount;  
32.     //确保线程数足够  
33.     histo_kernel << <blocks * 2, (SIZE + 2 * blocks - 1) / blocks / 2 >> > (SIZE, dev_histo);  
34.   
35.     //数据拷贝回CPU内存  
36.     cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);  
37.     printf("Threads SUM：%d\n", threadSum);  
38.     getchar();  
39.     cudaFree(dev_histo);  
40.     return 0;  
41. }  

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