交換運算子的單塊並行約簡

在 CUDA 中並行縮減的最簡單方法是分配一個塊來執行任務：

static const int arraySize = 10000;
static const int blockSize = 1024;

__global__ void sumCommSingleBlock(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;
    __syncthreads();
    for (int size = blockSize/2; size>0; size/=2) { //uniform
        if (idx<size)
            r[idx] += r[idx+size];
        __syncthreads();
    }
    if (idx == 0)
        *out = r[0];
}

...

sumCommSingleBlock<<<1, blockSize>>>(dev_a, dev_out);

當資料大小不是很大（幾個 thousants 元素）時，這是最可行的。當減少是一些更大的 CUDA 計劃的一部分時，通常會發生這種情況。如果輸入從一開始就匹配 blockSize，則可以完全刪除第一個 for 迴圈。

請注意，在第一步中，當元素多於執行緒時，我們完全獨立地新增內容。只有當問題減少到 blockSize 時，實際的並行減少才會觸發。相同的程式碼可以應用於任何其他可交換的關聯運算子，例如乘法，最小值，最大值等。

注意，可以使演算法更快，例如通過使用扭曲級並行縮減。