僅使用暫存器進行單經線並行縮減

通常，減少是在全域性或共享陣列上執行的。但是，當縮小是在非常小的範圍內執行時，作為更大的 CUDA 核心的一部分，它可以用單個 warp 執行。當發生這種情況時，在 Keppler 或更高版本的架構上（CC> = 3.0），可以使用 warp-shuffle 函式來避免使用共享記憶體。

例如，假設 warp 中的每個執行緒都包含單個輸入資料值。所有執行緒一起有 32 個元素，我們需要總結（或執行其他關聯操作）

__device__ int sumSingleWarpReg(int value) {
    value += __shfl_down(value, 1);
    value += __shfl_down(value, 2);
    value += __shfl_down(value, 4);
    value += __shfl_down(value, 8);
    value += __shfl_down(value, 16);
    return __shfl(value,0);
}

此版本適用於可交換和非交換運算子。