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

#include

#define size 1024

__global__ void

histo_kernel(

intsize, unsigned

int*histo)

} intmain(

void

)

使用原子操作正確的結果是523776，不使用原子操作的結果不確定，其中一次執行結果是711，顯然是不對的。

CUDA atomic原子操作

CUDA atomic原子操作

5 1 CUDA atomic原子操作

原子性，原子操作

CUDA atomic原子操作

CUDA atomic原子操作

5 1 CUDA atomic原子操作

原子性，原子操作

相關推薦