No 7 5 OpenCL 同步原子操作

當多個執行緒在不同處理器上執行，同時訪問相同記憶體時會存在競態條件，使用原子操作可以避免競爭。原子操作往往會對匯流排做一次鎖步操作（lock-step），讓當前匯流排上的訪存操作能按照次序執行。同時又會重新整理cache，使得任一線程對全域性變數使用了原子操作之後，其它所有執行緒都可見。

使用原子操作做同步開銷是相當大的，但是相對於使用更原始的阻塞當前執行緒執行的同步方式而言又是比較高效的。因此，當對某些特定資料做同步更新時，不需要使用柵欄（fence）等這種更低效的同步處理機制，我們可以直接對那些儲存位址採用原子操作。

opencl 提供了下面一系列原子操作。

int atomic_add (volatile __global int *p, int val)

在乙個原子事務中執行。讀取 p 指向位置的內容（用作返回值），將 p 指向位置的內容加上 val 後再存入該位置。

核心在多個裝置上對同一記憶體位置執行的原子操作沒法保證原子性。

原子減法操作atomic_sub和加法操作類似。

int atomic_xchg (volatile __global int *p, int val)

原子交換操作常用於對某個變數進行初始化或是作為互斥體（mutex）使用。將 val 存入 p 指向的位置，返回 p 位置修改之前的內容。

int atomic_inc(volatile __global int *p)

原子加 1 操作。讀取 p 指向位置的內容（用作返回值），將 p 指向位置的內容加上常量值1後再存入該位置。原子減 1 操作atomic_dec和加 1 操作類似。

int atomic_cmpxchg(volatile __global int *p, int cmp, int val)

如果 p 指向位置的內容和 cmp 相等，則將 val 存放 p 指向的位置，否則 p 指向位置的內容不變。返回 p 位置修改前的內容。

int atomic_min (volatile __local int *p, int val)

以原子的方式求最小值。讀取 p 指向位置的內容（用作返回值），並將該值和val比較，然後在 p 指向的位置存入較小的值。求最大值atomic_max操作和求最小值操作類似。

int atomic_and (volatile __global int *p, int val)

以原子的方式執行按位與運算操作。讀取 p 指向位置的內容（用作返回值），並將讀取的值和val執行與操作，然後將將結果存入 p 指向的位置。原子按位或操作atomic_or和原子按位異或atomic_xor操作與原子按位與操作類似。

示例程式 openclsyncatomic 以工作組為單位，分別對工作組中每個工作項做點乘運算，然後將運算結果以原子的方式相加到全域性記憶體變數dst中，供主機端獲取。

__kernel void kernel_dot(__global int *dst, __global int *src1,  __global int *src2)
}

No 7 5 OpenCL 同步 原子操作