CUDA中的Warp Shuffle

Warp Shuffle Functions

__shfl_sync、__shfl_up_sync、__shfl_down_sync 和 __shfl_xor_sync 在 warp 內的線程之間交換變量。

由計算能力 3.x 或更高版本的設備支持。

棄用通知：__shfl、__shfl_up、__shfl_down 和 __shfl_xor 在 CUDA 9.0 中已針對所有設備棄用。

刪除通知：當面向具有 7.x 或更高計算能力的設備時，__shfl、__shfl_up、__shfl_down 和 __shfl_xor 不再可用，而應使用它們的同步變體。

作者添加:這里可能大家對接下來會提到的threadIndex, warpIdx, laneIndex會比較混淆.那么我用下圖來說明.

1. Synopsis

T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize); T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize); T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width=warpSize); T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);

T 可以是 int、unsigned int、long、unsigned long、long long、unsigned long long、float 或 double。 包含 cuda_fp16.h 頭文件后，T 也可以是 __half 或 __half2。 同樣，包含 cuda_bf16.h 頭文件后，T 也可以是 __nv_bfloat16 或 __nv_bfloat162。

2. Description

__shfl_sync() 內在函數允許在 warp 內的線程之間交換變量，而無需使用共享內存。 交換同時發生在 warp 中的所有活動線程（并以mask命名），根據類型移動每個線程 4 或 8 個字節的數據。

warp 中的線程稱為通道(lanes)，并且可能具有介于 0 和 warpSize-1（包括）之間的索引。 支持四種源通道(source-lane)尋址模式：

__shfl_sync()

從索引通道直接復制

__shfl_up_sync()

從相對于調用者 ID 較低的通道復制

__shfl_down_sync()

從相對于調用者具有更高 ID 的通道復制

__shfl_xor_sync()

基于自身通道 ID 的按位異或從通道復制

線程只能從積極參與 __shfl_sync() 命令的另一個線程讀取數據。 如果目標線程處于非活動狀態，則檢索到的值未定義。

所有 __shfl_sync() 內在函數都采用一個可選的寬度參數，該參數會改變內在函數的行為。 width 的值必須是 2 的冪； 如果 width 不是 2 的冪，或者是大于 warpSize 的數字，則結果未定義。

__shfl_sync() 返回由 srcLane 給定 ID 的線程持有的 var 的值。 如果 width 小于 warpSize，則 warp 的每個子部分都表現為一個單獨的實體，其起始邏輯通道 ID 為 0。如果 srcLane 超出范圍 [0:width-1]，則返回的值對應于通過 srcLane srcLane modulo width所持有的 var 的值 （即在同一部分內）。

作者添加:這里原本中說的有點繞,我還是用圖來說明比較好.注意下面四個圖均由作者制作,如果有問題,僅僅是作者水平問題-_-!.

__shfl_up_sync() 通過從調用者的通道 ID 中減去 delta 來計算源通道 ID。 返回由生成的通道 ID 保存的 var 的值：實際上， var 通過 delta 通道向上移動。 如果寬度小于 warpSize，則warp的每個子部分都表現為一個單獨的實體，起始邏輯通道 ID 為 0。源通道索引不會環繞寬度值，因此實際上較低的 delta 通道將保持不變。

__shfl_down_sync() 通過將 delta 加調用者的通道 ID 來計算源通道 ID。 返回由生成的通道 ID 保存的 var 的值：這具有將 var 向下移動 delta 通道的效果。 如果 width 小于 warpSize，則 warp 的每個子部分都表現為一個單獨的實體，起始邏輯通道 ID 為 0。至于 __shfl_up_sync()，源通道的 ID 號不會環繞寬度值，因此 upper delta lanes將保持不變。

__shfl_xor_sync() 通過對調用者的通道 ID 與 laneMask 執行按位異或來計算源通道 ID：返回結果通道 ID 所持有的 var 的值。 如果寬度小于warpSize，那么每組寬度連續的線程都能夠訪問早期線程組中的元素，但是如果它們嘗試訪問后面線程組中的元素，則將返回他們自己的var值。 這種模式實現了一種蝶式尋址模式，例如用于樹規約和廣播。

新的 *_sync shfl 內部函數采用一個掩碼，指示參與調用的線程。 必須為每個參與線程設置一個表示線程通道 ID 的位，以確保它們在硬件執行內部函數之前正確收斂。 掩碼中命名的所有非退出線程必須使用相同的掩碼執行相同的內在函數，否則結果未定義。

3. Notes

線程只能從積極參與 __shfl_sync() 命令的另一個線程讀取數據。 如果目標線程處于非活動狀態，則檢索到的值未定義。

寬度必須是 2 的冪（即 2、4、8、16 或 32）。 未指定其他值的結果。

4. Examples

4.1. Broadcast of a single value across a warp

#include <stdio.h>__global__ void bcast(int arg) {int laneId = threadIdx.x & 0x1f;int value;if (laneId == 0) // Note unused variable forvalue = arg; // all threads except lane 0value = __shfl_sync(0xffffffff, value, 0); // Synchronize all threads in warp, and get "value" from lane 0if (value != arg)printf("Thread %d failed.\n", threadIdx.x); }int main() {bcast<<< 1, 32 >>>(1234);cudaDeviceSynchronize();return 0; }

4.2. Inclusive plus-scan across sub-partitions of 8 threads

#include <stdio.h>__global__ void scan4() {int laneId = threadIdx.x & 0x1f;// Seed sample starting value (inverse of lane ID)int value = 31 - laneId;// Loop to accumulate scan within my partition.// Scan requires log2(n) == 3 steps for 8 threads// It works by an accumulated sum up the warp// by 1, 2, 4, 8 etc. steps.for (int i=1; i<=4; i*=2) {// We do the __shfl_sync unconditionally so that we// can read even from threads which won't do a// sum, and then conditionally assign the result.int n = __shfl_up_sync(0xffffffff, value, i, 8);if ((laneId & 7) >= i)value += n;}printf("Thread %d final value = %d\n", threadIdx.x, value); }int main() {scan4<<< 1, 32 >>>();cudaDeviceSynchronize();return 0; }

4.3. Reduction across a warp

#include <stdio.h>__global__ void warpReduce() {int laneId = threadIdx.x & 0x1f;// Seed starting value as inverse lane IDint value = 31 - laneId;// Use XOR mode to perform butterfly reductionfor (int i=16; i>=1; i/=2)value += __shfl_xor_sync(0xffffffff, value, i, 32);// "value" now contains the sum across all threadsprintf("Thread %d final value = %d\n", threadIdx.x, value); }int main() {warpReduce<<< 1, 32 >>>();cudaDeviceSynchronize();return 0; }

總結

以上是生活随笔為你收集整理的CUDA中的Warp Shuffle的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。