非交換運算元的單經線並行約簡

有時，減少必須在非常小的範圍內執行，作為更大的 CUDA 核心的一部分。例如，假設輸入資料恰好有 32 個元素 - warp 中的執行緒數。在這種情況下，可以分配單個扭曲來執行減少。鑑於 warp 在完美同步中執行，與塊級減少相比，可以刪除許多 __syncthreads() 指令。

static const int warpSize = 32;

__device__ int sumNoncommSingleWarp(volatile int* shArr) {
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    if (idx%2 == 0) shArr[idx] += shArr[idx+1];
    if (idx%4 == 0) shArr[idx] += shArr[idx+2];
    if (idx%8 == 0) shArr[idx] += shArr[idx+4];
    if (idx%16 == 0) shArr[idx] += shArr[idx+8];
    if (idx == 0) shArr[idx] += shArr[idx+16];
    return shArr[0];
}

shArr 最好是共享儲存器中的陣列。warp 中的所有執行緒的值應該相同。如果 sumCommSingleWarp 被多個 warp 呼叫，則 shArr 應該在 warp 之間不同（每個 warp 中相同）。

引數 shArr 被標記為 volatile，以確保在指示的位置實際執行對陣列的操作。否則，對 shArr[idx] 的重複分配可以優化為對暫存器的分配，只有最終分配是實際儲存到 shArr。當發生這種情況時，其他執行緒看不到即時分配，從而產生不正確的結果。注意，你可以將普通的非易失性陣列作為 volatile 引數傳遞，與將非 const 作為 const 引數傳遞時相同。

如果一個人不關心 shArr[1..31] 的最終內容並且可以用零填充 shArr[32..47]，可以簡化上面的程式碼：

static const int warpSize = 32;

__device__ int sumNoncommSingleWarpPadded(volatile int* shArr) {
    //shArr[32..47] == 0
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    shArr[idx] += shArr[idx+1];
    shArr[idx] += shArr[idx+2];
    shArr[idx] += shArr[idx+4];
    shArr[idx] += shArr[idx+8];
    shArr[idx] += shArr[idx+16];
    return shArr[0];
}

在此設定中，我們刪除了所有 if 條件，這些條件構成了大約一半的指令。額外的執行緒執行一些不必要的新增，將結果儲存到 shArr 的單元格中，最終對最終結果沒有影響。由於 warp 在 SIMD 模式下執行，我們實際上並沒有通過讓這些執行緒無所事事來節省時間。