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，如果不加原子操作，執行結果是錯誤並且是不確定的。

1. #include <stdio.h>    
2. #include <stdlib.h>   
3. #include <cuda_runtime.h>  
4. #define SIZE 1024
5. __global__ void histo_kernel(int size, unsigned int *histo)  
6. {  
7.     int i = threadIdx.x + blockIdx.x * blockDim.x;  
8.     if (i < size)  
9.     {  
10.         //*histo+=i;
11.         atomicAdd(histo, i);  
12.     }  
13. }  
14. int main(void)  
15. {  
16.     int threadSum = 0;  
17.     //分配記憶體並拷貝初始資料
18.     unsigned int *dev_histo;  
19.     cudaMalloc((void**)&dev_histo, sizeof(int));  
20.     cudaMemcpy(dev_histo, &threadSum, sizeof(int), cudaMemcpyHostToDevice);  
21.     // kernel launch - 2x the number of mps gave best timing  
22.     cudaDeviceProp  prop;  
23.     cudaGetDeviceProperties(&prop, 0);  
24.     int blocks = prop.multiProcessorCount;  
25.     //確保執行緒數足夠
26.     histo_kernel << <blocks * 2, (SIZE + 2 * blocks - 1) / blocks / 2 >> > (SIZE, dev_histo);  
27.     //資料拷貝回CPU記憶體
28.     cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);  
29.     printf("Threads SUM：%d\n", threadSum);  
30.     getchar();  
31.     cudaFree(dev_histo);  
32.     return 0;  
33. }