非交換運算元的單塊並行約簡
與可交換版本相比,對非交換運算子進行並行縮減更為複雜。在示例中,為簡單起見,我們仍然使用整數加法。例如,它可以用矩陣乘法代替,它實際上是非交換的。注意,這樣做時,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 級別和刪除共享記憶體庫衝突。