交换算子的单经线并联约简

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

static const int warpSize = 32;

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

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

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

如果在减少之后不关心 shArr[1..31] 的内容，可以进一步简化代码：

static const int warpSize = 32;

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

在此设置中，我们删除了许多 if 条件。额外的线程执行一些不必要的添加，但我们不再关心它们产生的内容。由于 warp 在 SIMD 模式下执行，我们实际上并没有通过让这些线程无所事事来节省时间。另一方面，评估条件确实需要相对较长的时间，因为这些条款的主体非常小。如果 shArr[32..47] 填充 0，则可以删除初始 if 语句。

经线级减少也可用于提高块级减少：

__global__ void sumCommSingleBlockWithWarps(const int *a, int *out) {
    int idx = threadIdx.x;
    int sum = 0;
    for (int i = idx; i < arraySize; i += blockSize)
        sum += a[i];
    __shared__ int r[blockSize];
    r[idx] = sum;
    sumCommSingleWarp(&r[idx & ~(warpSize-1)]);
    __syncthreads();
    if (idx<warpSize) { //first warp only
        r[idx] = idx*warpSize<blockSize ? r[idx*warpSize] : 0;
        sumCommSingleWarp(r);
        if (idx == 0)
            *out = r[0];
    }
}

论坛 &r[idx & ~(warpSize-1)] 基本上是 r + warpIdx*32。这有效地将 r 阵列分成 32 个元素的块，并且每个块被分配给单独的 warp。