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

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

static const int warpSize = 32;

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

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

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

如果在減少之後不關心 shArr[1..31] 的內容，可以進一步簡化程式碼：

static const int warpSize = 32;

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

在此設定中，我們刪除了許多 if 條件。額外的執行緒執行一些不必要的新增，但我們不再關心它們產生的內容。由於 warp 在 SIMD 模式下執行，我們實際上並沒有通過讓這些執行緒無所事事來節省時間。另一方面，評估條件確實需要相對較長的時間，因為這些條款的主體非常小。如果 shArr[32..47] 填充 0，則可以刪除初始 if 語句。

經線級減少也可用於提高塊級減少：

__global__ void sumCommSingleBlockWithWarps(const int *a, int *out) {
    int idx = threadIdx.x;
    int sum = 0;
    for (int i = idx; i < arraySize; i += blockSize)
        sum += a[i];
    __shared__ int r[blockSize];
    r[idx] = sum;
    sumCommSingleWarp(&r[idx & ~(warpSize-1)]);
    __syncthreads();
    if (idx<warpSize) { //first warp only
        r[idx] = idx*warpSize<blockSize ? r[idx*warpSize] : 0;
        sumCommSingleWarp(r);
        if (idx == 0)
            *out = r[0];
    }
}

論壇 &r[idx & ~(warpSize-1)] 基本上是 r + warpIdx*32。這有效地將 r 陣列分成 32 個元素的塊，並且每個塊被分配給單獨的 warp。