CUDA 规约优化实现

归约（Reduction）是并行计算中的经典操作，例如对数组求和。在 CUDA 中，由于 GPU 的 SIMT 架构，实现高效的归约需要充分考虑线程协作、内存层次和同步开销。本文将基于一份示例代码，逐步实现 CUDA 归约的优化。

问题：把 N=256M 个 float 加在一起。理想带宽利用 = N×4 字节 / 时间 / peak。

一、实现一：divergent 版

for (int s = 1; s < blockDim.x; s *= 2) {
    if (tid % (2*s) == 0) sdata[tid] += sdata[tid + s];   // ❌ warp 内 31/32 lane 闲着
    __syncthreads();
}

判断条件 tid % (2*s) == 0 让 warp 内每次只活 1/2、1/4 … 的 lane → warp divergence。

二、实现二：顺序地址

for (int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (tid < s) sdata[tid] += sdata[tid + s];     // 同 warp 内 tid 都满足条件
    __syncthreads();
}

每 warp 全员活动直到 s = 16，分歧大幅减少；同时访问也 coalesce。

三、实现三：启动减半 + 每 thread 加两元素

int i = blockIdx.x * (blockDim.x * 2) + tid;
float v = in[i] + in[i + blockDim.x];      // load 两个再写 sdata
sdata[tid] = v;

kernel launch 开销摊薄一半；同时让算术与内存比例上升。

四、实现四：展开最后一个 warp（去掉无用 sync）

for (int s = blockDim.x / 2; s > 32; s >>= 1) {
    if (tid < s) sdata[tid] += sdata[tid + s];
    __syncthreads();
}
// 当 s ≤ 32, 整个折叠都在同一 warp 内，硬件已经 SIMT 同步，不需要 __syncthreads
if (tid < 32) warp_reduce(sdata, tid);

五、warp shuffle

// 用 __shfl_down_sync 直接在寄存器间交换
for (int off = 16; off > 0; off >>= 1)
    v += __shfl_down_sync(0xffffffff, v, off);
// 32 个 warp-sum 写 shared, 让 warp 0 再 reduce 一次得到 block-sum

性能采集表：

版本	关键改动	时间	带宽
v1	原始	0.117 ms	1.4 GB/s
v2	sequential addressing	0.045 ms	3.6 GB/s
v3	每 thread 加 2 个	0.038 ms	4.3 GB/s
v4	unroll last warp	0.035 ms	4.7 GB/s
v5	warp shuffle	0.031 ms	5.1 GB/s