CUDA 的塊間同步方法

cuda 塊間同步方法有以下三種

① Lock-Based Synchronization

② Lock-Free Synchronization

③ __threadfence(）

基于鎖的塊間同步

cuda 基于鎖的同步的基本思想是使用一個全局互斥量變量來計算到達同步點的線程塊的數量。如下代碼所示，在 barrier 函數 __gpu_sync() 中，在一個塊完成它的計算之后，它的一個線程 (這里人為設置為?0 號線程，我們稱之為主導線程) 將自動地向 g_mutex 添加 1 （原子操作）。然后，主導線程將重復將?g_mutex 和一個目標值 goalVal 進行比較。如果 g_mutex 等于 goalVal，那么就意味著同步完成，即每個線程塊都可以進行下一階段的計算。在該設計中，當第一次調用 barrier 函數時，將 goalVal 設置為內核中的塊數 N 。然后，當連續調用 barrier 函數時，goalVal的值每次遞增 N 。這種設計比保持 goalVal 常量并在每個 barrier 之后重新設置 g_mutex 更有效，因為前者節省了指令的數量并避免了條件分支 。

// lock-based

__device__ volatile int g_mutex;

// GPU lock-based synchronization function

__device__ void __gpu_sync(int goalVal)

{

// thread ID in a block

int tid_in_block = threadIdx.x * blockDim.y + threadIdx.y;

// only thread 0 is used for synchronization

if (tid_in_block == 0)

{

atomicAdd((int*) &g_mutex, 1);

// only when all blocks add 1 go g_mutex

// will g_mutex equal to goalVal

while (g_mutex != goalVal)

{

// Do nothing here

}

}

__syncthreads();

}

無鎖的塊間同步

在 CUDA 基于鎖的同步中，互斥量變量 g_mutex 是通過原子函數 atomicAdd() 添加的。這意味著 g_mutex 的添加只能按順序執行，即使這些操作是在不同的塊中執行的。因此，提出一種完全避免使用原子操作的無鎖同步方法。這種方法的基本思想是為每個線程塊分配一個同步變量，這樣每個塊就可以獨立地記錄其同步狀態，而不必爭用單個全局互斥鎖變量。如下代碼所示，我們的無鎖同步方法使用兩個數組 Arrayin 和 Arrayout 來協調來自不同塊的同步請求。在這兩個數組中，每個元素都映射到內核中的一個線程塊，即，將元素 i 映射到線程塊 i ，算法大致分為三個步驟:

一開始，當 block i 準備好通信時，它的主導線程 (線程 0) 將 Arrayin 中的元素 i 設置為目標值 goalVal。block i 中的主導線程然后忙等 Arrayout 的元素 i 被設置為 goalVal 。

然后人為的讓 block 1中的前 N 個（N 等于塊數）線程重復檢查 Arrayin 中的所有元素是否等于 goalVal ，線程 i 負責檢查 Arrayin 中的第 i 個元素（一對一檢查）。將 Arrayin 中的所有元素設置為 goalVal 后，每個檢查線程將 Arrayout 中的相應元素設置為 goalVal 。注意，在更新 Arrayout的元素之前，block 1 的每個檢查線程都會調用塊內 barrier 函數? __syncthreads()。

最后，當 block i 的主導線程看到 Arrayout 中的對應元素被設置為 goalVal 時，該 block 將繼續執行。

// lock-free

__device__ void __gpu_sync(int goalVal, volatile int *Arrayin, volatile int *Arrayout)

{

// thread ID in a block

int tid_in_blk = threadIdx.x * blockDim.y + threadIdx.y;

int nBlockNum = gridDim.x * gridDim.y;

int bid = blockIdx.x * gridDim.y + blockIdx.y;

// only thread 0 is used for synchonization

if (tid_in_blk == 0)

{

Arrayin[bid] = goalVal;

}

if (bid == 1)

{

if (tid_in_blk < nBlockNum)

{

while (Arrayin[tid_in_blk] != goalVal)

{

// Do nothing here

}

}

__syncthreads();

if (tid_in_blk < nBlockNum)

{

Arrayout[tid_in_blk] = goalVal;

}

}

if (tid_in_blk = 0)

{

while (Arrayout[bid] != goalVal)

{

// Do nothing here

}

}

__syncthreads();

}

從以上代碼可以看出，CUDA 無鎖同步中沒有原子操作。所有的操作都可以并行執行。不同線程塊的同步由單個塊 (block?1) 中的 N 個線程來控制，可以通過調用塊內 barrier 函數 __syncthreads() 來有效地同步。

__threadfence()

最后，值得注意的是，另外一種保證?CUDA 塊間同步通信的正確性的辦法是使用?__threadfence() （CUDA 2.2中引入了一個新的函數 ）。這個函數將阻塞調用線程，直到之前對 全局內存 或 共享內存 的寫入對其他線程可見為止。但是使用 __threadfence() 也會引起一定的額外開銷，所以需要進行實際測試和權衡。

任務調度

版權聲明：本文內容由網絡用戶投稿，版權歸原作者所有，本站不擁有其著作權，亦不承擔相應法律責任。如果您發現本站中有涉嫌抄襲或描述失實的內容，請聯系我們jiasou666@gmail.com 處理，核實后本網站將在24小時內刪除侵權內容。

版權聲明：本文內容由網絡用戶投稿，版權歸原作者所有，本站不擁有其著作權，亦不承擔相應法律責任。如果您發現本站中有涉嫌抄襲或描述失實的內容，請聯系我們jiasou666@gmail.com 處理，核實后本網站將在24小時內刪除侵權內容。