仅使用寄存器进行单经线并行缩减

通常，减少是在全局或共享阵列上执行的。但是，当缩小是在非常小的范围内执行时，作为更大的 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);
}

此版本适用于可交换和非交换运算符。