非交换算子的单经线并行约简

有时，减少必须在非常小的范围内执行，作为更大的 CUDA 内核的一部分。例如，假设输入数据恰好有 32 个元素 - warp 中的线程数。在这种情况下，可以分配单个扭曲来执行减少。鉴于 warp 在完美同步中执行，与块级减少相比，可以删除许多 __syncthreads() 指令。

static const int warpSize = 32;

__device__ int sumNoncommSingleWarp(volatile int* shArr) {
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    if (idx%2 == 0) shArr[idx] += shArr[idx+1];
    if (idx%4 == 0) shArr[idx] += shArr[idx+2];
    if (idx%8 == 0) shArr[idx] += shArr[idx+4];
    if (idx%16 == 0) shArr[idx] += shArr[idx+8];
    if (idx == 0) shArr[idx] += shArr[idx+16];
    return shArr[0];
}

shArr 最好是共享存储器中的数组。warp 中的所有线程的值应该相同。如果 sumCommSingleWarp 被多个 warp 调用，则 shArr 应该在 warp 之间不同（每个 warp 中相同）。

参数 shArr 被标记为 volatile，以确保在指示的位置实际执行对阵列的操作。否则，对 shArr[idx] 的重复分配可以优化为对寄存器的分配，只有最终分配是实际存储到 shArr。当发生这种情况时，其他线程看不到即时分配，从而产生不正确的结果。注意，你可以将普通的非易失性数组作为 volatile 参数传递，与将非 const 作为 const 参数传递时相同。

如果一个人不关心 shArr[1..31] 的最终内容并且可以用零填充 shArr[32..47]，可以简化上面的代码：

static const int warpSize = 32;

__device__ int sumNoncommSingleWarpPadded(volatile int* shArr) {
    //shArr[32..47] == 0
    int idx = threadIdx.x % warpSize; //the lane index in the warp
    shArr[idx] += shArr[idx+1];
    shArr[idx] += shArr[idx+2];
    shArr[idx] += shArr[idx+4];
    shArr[idx] += shArr[idx+8];
    shArr[idx] += shArr[idx+16];
    return shArr[0];
}

在此设置中，我们删除了所有 if 条件，这些条件构成了大约一半的指令。额外的线程执行一些不必要的添加，将结果存储到 shArr 的单元格中，最终对最终结果没有影响。由于 warp 在 SIMD 模式下执行，我们实际上并没有通过让这些线程无所事事来节省时间。