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

與可交換版本相比,對非交換運算子進行並行縮減更為複雜。在示例中,為簡單起見,我們仍然使用整數加法。例如,它可以用矩陣乘法代替,它實際上是非交換的。注意,這樣做時,0 應該被乘法的中性元素替換 - 即單位矩陣。

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

__global__ void sumNoncommSingleBlock(const int *gArr, int *out) {
    int thIdx = threadIdx.x;
    __shared__ int shArr[blockSize*2];
    __shared__ int offset;
    shArr[thIdx] = thIdx<arraySize ? gArr[thIdx] : 0;
    if (thIdx == 0)
        offset = blockSize;
    __syncthreads();
    while (offset < arraySize) { //uniform
        shArr[thIdx + blockSize] = thIdx+offset<arraySize ? gArr[thIdx+offset] : 0;
        __syncthreads();
        if (thIdx == 0)
            offset += blockSize;
        int sum = shArr[2*thIdx] + shArr[2*thIdx+1];
        __syncthreads();
        shArr[thIdx] = sum;
    }
    __syncthreads();
    for (int stride = 1; stride<blockSize; stride*=2) { //uniform
        int arrIdx = thIdx*stride*2;
        if (arrIdx+stride<blockSize)
            shArr[arrIdx] += shArr[arrIdx+stride];
        __syncthreads();
    }   
    if (thIdx == 0)
        *out = shArr[0];
}

...

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

在第一個 while 迴圈中執行,只要有多個輸入元素而不是執行緒。在每次迭代中,執行單次縮減,並將結果壓縮到 shArr 陣列的前半部分。然後下半部分將填充新資料。

gArr 載入所有資料後,執行第二個迴圈。現在,我們不再壓縮結果(這需要額外費用 3)。在每個步驟中,執行緒 n 訪問 2*n-th 活動元素並將其與 tihuan 第 5 個元素相加:

有許多方法可以進一步優化這個簡單的例子,例如通過減少 warp 級別和刪除共享記憶體庫衝突。