CUDA專業小貼士：通過 Warp-聚合的原子操作來優化過濾

注：這篇文章已經（于 2017年11月）更新為 CUDA 9 和最新的 GPU。NVCC編譯器現在可以在許多情況下自動為原子操作執行 warp-聚合，因此無需額外的工作就可以獲得更高的性能。事實上，編譯器生成的代碼比手工編寫的 warp 聚合代碼更快。這篇文章主要是為那些想了解它是如何工作的，并將類似的技術應用于其他問題的人準備的。

在本文中，我將介紹 warp-聚合的原子操作，這是一種可以在多個線程原子地累加到單個計數器時提高性能的有用技術。在 warp-聚合中，warp 的線程首先計算它們之間的總增量，然后選擇單個線程將增量原子地添加到全局計數器中。這種聚合減少了原子操作的數量【最多可以是一個 warp 中的線程數量(當前gpu上最多可以減少 32x )】，并且可以顯著提高性能。此外，在許多典型情況下，可以將標準原子操作順便替換為 warp-聚合實現，因此它可以作為提高復雜應用程序性能的一種簡單方法。

問題: 通過斷定過濾

考慮如下過濾問題：我有一個源數組 src，包含 n 個元素和一個斷定，我需要將滿足斷定的 src 的所有元素復制到目標數組 dst 中。為了簡單起見，假設 dst 的長度至少為 n，并且 dst 數組中元素的順序無關緊要。對于這個例子，我假設數組元素是整數，并且只有當元素為正時判定才為真。下面是過濾的 CPU 實現示例。

int filter(int *dst, const int *src, int n) {

int nres = 0;

for (int i = 0; i < n; i++)

if (src[i] > 0)

dst[nres++] = src[i];

// return the number of elements copied

return nres;

}

過濾(也稱為流壓縮)是一種常見的操作，它是許多編程語言的標準庫的一部分，可以使用多種名稱，包括 grep、copy_if、select、FindAll 等等。它也經常被簡單地實現為一個循環，因為它可能與周圍的代碼緊密集成。

結合全局和共享內存的解決方案

現在，如果我想在? GPU 上實現過濾，且并行處理數組 src 的元素，該怎么辦? 一種直接的方法是使用一個全局計數器，并對 dst 數組中寫入的每個新元素原子地遞增它。這個的 GPU 實現可能如下所示。

__global__

void filter_k(int *dst, int *nres, const int *src, int n) {

int i = threadIdx.x + blockIdx.x * blockDim.x;

if(i < n && src[i] > 0)

dst[atomicAdd(nres, 1)] = src[i];

}

這個實現的主要問題是，從 src 讀取正元素的（網格中的）所有線程都會累加一個計數器 nres 。根據正元素的數量，這可能是非常多的線程。因此，atomicAdd()?的沖突程度很高，這限制了性能。您可以在 圖1 中看到這一點，它繪制了 Kepler K80 GPU在處理1億個(100*2^20)元素時獲得的內核帶寬(包括讀和寫，但不包括原子操作)。

圖1.? 開普勒 K80 GPU 上的全局原子操作濾波的性能 (CUDA 8.0.61) 。

帶寬與執行的原子操作數或數組中正元素的比例成反比。對于5%的部分（fraction），性能是可以接受的(大約55 GiB/s)，但是當更多的元素通過過濾器時，性能會急劇下降，對于 50% 的部分 (fraction)，性能僅為 8 GiB/s左右。原子操作顯然是一個瓶頸，需要刪除或減少原子操作以提高應用程序性能。

提高過濾性能的一種方法是使用共享內存執行原子操作。這提高了每個操作的速度，并減少了沖突的程度，因為計數器只在單個塊中的線程之間共享。使用這種方法，每個線程塊只需要一個全局 atomicAdd()。下面是用這種方法實現的內核。

__global__

void filter_shared_k(int *dst, int *nres, const int* src, int n) {

__shared__ int l_n;

int i = blockIdx.x * (NPER_THREAD * BS) + threadIdx.x;

for (int iter = 0; iter < NPER_THREAD; iter++) { // 迭代 NPER_THREAD 次

// zero the counter

if (threadIdx.x == 0)

l_n = 0;

__syncthreads();

// get the value, evaluate the predicate, and

// increment the counter if needed

int d, pos;

if(i < n) {

d = src[i];

if(d > 0)

pos = atomicAdd(&l_n, 1); // 在共享內存上進行原子加（l_n 不會超過 BS）

}

__syncthreads();

// leader increments the global counter

if(threadIdx.x == 0)

l_n = atomicAdd(nres, l_n); // 注意，這邊是用 nres 的舊值來覆蓋 l_n

__syncthreads();

// threads with true predicates write their elements

if(i < n && d > 0) {

pos += l_n; // increment local pos by global counter 下一塊（塊內）各個線程的位置

dst[pos] = d;

}

__syncthreads();

i += BS; // 一次迭代跳過 BS 個 src (一個線程塊有 BS 個線程)

}

}

另一種方法是首先使用一個并行前綴和來計算每個元素的輸出索引。Thrust 庫的 copy_if() 函數使用了這種方法的優化版本。圖2展示了開普勒? K80 的兩種方法的性能。盡管共享內存原子技術提高了過濾性能，但其性能仍然保持在原始方法的1.5倍以內。原子操作仍然是一個瓶頸，因為操作的數量沒有改變。Thurst 比兩種方法都適用于高過濾部分 (fraction)，但前期成本較大，不能用于小過濾部分的攤銷。

需要注意的是，與 Thrust 的比較不是嚴格的同類比較，因為? Thrust 實現了一個穩定的過濾器:它保留了輸出中輸入元素的相對順序。這是使用前綴和來實現它的結果，但其代價更高。如果我們不需要一個穩定的過濾器，那么純原子方法更簡單，執行的工作也更少。

圖2.? 基于Kepler K80 GPU 的共享內存原子操作濾波的性能(CUDA 8.0.61)。

Warp-聚合的原子操作

warp-聚合是將來自一個 warp 中的多個線程的原子操作組合成一個原子操作的過程。這種方法與使用共享內存是正交的（完全不同的）：原子的類型保持不變，但是我們使用的原子操作更少。使用 warp-聚合，我們使用以下步驟替換原子操作。

從 warp 中的線程中選擇一個領導線程。

warp 中的所有線程都計算 warp 的總原子增量。

領導線程執行一次原子加法來計算 warp 的偏移量。

領導線程將偏移量廣播給 warp 中的所有其他線程。

每個線程都將自己在 warp 中的索引加上 warp 偏移量，以獲得其在輸出數組中的位置。

從 CUDA 9.0 開始，有兩個 api 可用來實現這一點：

協作組，一種用于管理協作線程組的 CUDA 編程模型的擴展；

warp 同步原語函數。

在執行一個 warp-聚合的原子操作之后，每個線程按照原始代碼的方式繼續，并將其值寫入 dst 數組中對應的位置。現在讓我們詳細考慮每個步驟。

步驟1：領導線程的選擇

在過濾中，可以重新組織代碼，使所有線程都處于活躍狀態。然而，在其他一些情況下，原子操作可能發生在嵌套條件中，其中一些線程可能處于非活躍狀態。通常，這種方法應該假設只有某些線程是活躍的，所以我需要一個由所有活躍線程組成的組。

要使用協作組，請包含頭文件并使用 cooperative_groups 命名空間。

#include

using namespace cooperative_groups;

將當前所有合并訪問的線程創建為一組。

auto g = coalesced_threads();

使用協作組能夠很容易獲得線程級別：調用 g.thread_rank() 。級別為 0 的線程將成為 leader。

如果您喜歡使用原語函數，可以從 _activemask() 開始。

unsigned int active = __activemask();

(一種較老的方法是使用 _ballot(1) 。這在 CUDA 8 上是可行的，但從 CUDA 9 開始就廢棄了。)

然后選出一個 leader 。warp 內的線程叫做 lane ;選擇 leader 最簡單的方法是使用數字最小的活躍 lane 。__ffs() 原語返回集合位（set bit）的基于 1 的最低索引，因此減去 1 得到基于 0 的索引。

int leader = __ffs(active) - 1;

步驟2: 計算總增量

對于過濾的示例，每個具有判定為真的線程將計數器遞增 1 。warp 的總增量等于活躍 lane 的數量(這里我不考慮不同 lane 增量不同的情況)。這對于協作組來說很簡單:?g.size()?返回組中的線程數。

如果您喜歡使用原語函數，您可以將由 _activemask() 返回的掩碼中 bits set 的個數作為總增量。為此，使用內置函數?_popc(int v) ，它返回整數 v 的二進制表示的 bits set 的個數。

int change = __popc(active);

步驟3：執行原子添加

只有領導線程(lane 0)執行原子操作。對于協作組，只需檢查 thread_rank() 是否返回0，就像這樣。

int warp_res;

if(g.thread_rank() == 0)

warp_res = atomicAdd(ctr, g.size());

如果喜歡使用原語函數，則必須使用?_lanemask_lt() 計算每個 lane 的 rank，該函數返回 ID小于當前 lane 的所有 lane (包括非活躍 lane)的掩碼。然后，您可以通過將這個掩碼與活躍的 lane 的掩碼進行與運算來計算 rank，并統計 bits set 的個數。

unsigned int rank = __popc(active & __lanemask_lt());

int warp_old;

if(rank == 0)

warp_old = atomicAdd(ctr, change); // ctr is the pointer to the counter

步驟4：廣播結果

在此步驟中，領導線程將 atomicAdd() 的結果廣播到 warp 中的其他 lane。我們可以通過在活躍 lane 上使用 shuffle 操作來實現這一點。

使用協作組，您可以使用 g.shfl(warp_res, 0) 廣播結果。?0 是領導線程的索引，它僅僅在活躍線程是組的一部分時才奏效(因為它是使用 coalesced_threads() 創建的)。

如果您喜歡使用原語函數，可以調用?_shfl_sync()，它具有以下簽名，其中 T 是32位或64位整數或浮點類型。

T __shfl_sync(unsigned int mask, T var, int srcLane, int width=warpSize);

shfl_sync() 返回由 srcLane 提供 ID 的線程所持有的值 var。mask 是參與調用的線程的掩碼。掩碼位為 1 的所有非退出線程（non-exited）必須使用相同的掩碼執行相同的內置函數，否則結果將是未定義的。 width 必須是 2 的冪次，且小于或等于 warp 尺寸。 warp 會按照該尺寸分成相同大小的組，srcLane 指的是組內的 lane 號。如果 srcLane 超出范圍[0:width-1](包括兩端)，則?srcLane 對 width 取模給出 lane 號。

下面的代碼使用了?_shfl_sync()?來廣播結果。

warp_res = __shfl_sync(active, warp_res, leader);

CUDA 8 和更早的實現使用了_shfl()，從 CUDA 9 開始就不提倡使用它 (已廢棄) 。

步驟5:計算每個 lane 的結果

最后一步計算每個 lane 的輸出位置，方法是將 warp 的廣播計數器的值添加到（活躍的 lanes?中的） lane 的 rank 上。

協作組的形式：

return g.shfl(warp_res, 0) + g.thread_rank();

原語函數的形式：

return warp_res + rank;

現在，我們可以將步驟 1-5 的代碼片連接起來，以獲得完整的 warp-聚合版本的增量函數。

對于協作組，代碼簡潔明了。

__device__ int atomicAggInc(int *ctr) {

auto g = coalesced_threads();

int warp_res;

if(g.thread_rank() == 0)

warp_res = atomicAdd(ctr, g.size());

return g.shfl(warp_res, 0) + g.thread_rank();

}

對于原語函數，代碼則更加復雜。

__device__ int atomicAggInc(int *ctr) {

unsigned int active = __activemask();

int leader = __ffs(active) - 1;

int change = __popc(active);

unsigned int rank = __popc(active & __lanemask_lt());

int warp_res;

if(rank == 0)

warp_res = atomicAdd(ctr, change);

warp_res = __shfl_sync(active, warp_res, leader);

return warp_res + rank;

}

性能對比

warp-聚合的原子增量函數是 atomicAdd(ctr, 1) 的一個替代，其中所有 warp 線程的 ctr 都是相同的。因此，我們可以使用atomicAggInc() 重寫 GPU 過濾，如下所示。

__global__ void filter_k(int *dst, const int *src, int n) {

int i = threadIdx.x + blockIdx.x * blockDim.x;

if(i >= n)

return;

if(src[i] > 0)

dst[atomicAggInc(nres)] = src[i];

}

注意，盡管我們在定義 warp-聚合時考慮了全局原子，但是沒有什么可以阻止對共享內存原子做同樣的事情。事實上，如果 ctr 是指向共享內存的指針，那么上面定義的 atomicAggInc(int *ctr) 函數就可以工作。因此，Warp-聚合還可以用于加速共享內存的過濾。圖3 顯示了有和沒有warp-聚合的不同類型的濾波在 開普勒 GPU 上性能的比較。

圖3.? 在Tesla K80(開普勒架構)GPU (CUDA 8.0.61)上不同過濾的性能。

對于開普勒 GPU 來說，帶有 warp-聚合的全局原子的版本顯然是贏家。它總是提供超過 80GiB?/s的帶寬，并且帶寬實際上隨著成功通過過濾器的元素的比例的升高而增加。這也表明原子不再是一個重要的瓶頸。與 全局原子 相比，性能提高了21倍。在同一個 GPU 上，一個簡單的復制操作的性能大約是190 GiB?/s。因此，我們可以說，使用 warp-聚合原子進行過濾的性能與簡單的復制操作相當。這也意味著過濾現在可以用于代碼的性能關鍵部分。還要注意，共享內存原子(帶有 warp-聚合)實際上比 warp-聚合原子慢。這表明 warp-聚合已經做得很好了，在開普勒上使用共享內存沒有帶來任何好處，只會帶來額外的開銷。

由于在某些情況下，可以使用 warp-聚合的 atomics 作為常規 atomics 的替代，所以編譯器現在在許多情況下自動執行這種優化就不足為奇了。事實上，編譯器從CUDA 7.5開始為后開普勒 GPU 做優化，而在CUDA 9中，它也為開普勒 GPU 做優化。因此，早期的比較是與開普勒上的 CUDA 8 進行的，在那里 warp 聚集的原子尚未自動插入。

圖4、圖5 和 圖6 顯示了開普勒、帕斯卡和伏特在?CUDA 9?上的性能比較，簡單 atomicAdd() 的性能類似于 warp-聚合的 atomics。

題圖4. 在 Tesla K80(Kepler) GPU 上不同過濾的性能(CUDA 9.0.176)

圖5. 在 Tesla P100 (Pascal) GPU (CUDA 9.0.176)上不同過濾的性能。

圖6. 在 Tesla V100 (Volta) GPU (CUDA 9.0.176)上不同過濾的性能。

總結

原子的 warp-聚合是一種有用的技術，可以提高在少量計數器上執行許多操作的應用程序的性能。在這篇文章中，我們將 warp-聚合應用于濾波，并在開普勒+CUDA 8 上獲得了一個量級以上的性能改進。事實上，這項技術非常有用，現在已經在 NVCC 編譯器中實現了，在很多情況下，默認情況下無需額外的工作就可以獲得 warp-聚合。

warp-聚合原子并不局限于過濾;您可以將它用于許多其他使用原子操作的應用程序中。

原文地址

任務調度 數據結構

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

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