CUDA并行归约优化：从Bank Conflict到Warp Shuffle

1. CUDA并行归约优化实战：从Bank Conflict到Warp Shuffle

作为一名CUDA开发者，并行归约(Parallel Reduction)是我们必须掌握的"Hello World"级算法。它看似简单——只是对一个数组求和，但却是理解GPU硬件特性的绝佳案例。本文将带你经历7个版本的迭代优化，从最初的2.3ms耗时逐步优化到0.6ms以下，深入剖析Bank Conflict、Warp Divergence等关键概念，最终利用Warp Shuffle指令达到性能极限。

2. 基础版本分析：V1与V2的演进

2.1 V1版本：交错寻址的陷阱

cpp复制template<int blockSize>
__global__ void reduce_v1(float *d_in, float *d_out){
    int tid = threadIdx.x;
    int gtid = threadIdx.x + blockIdx.x * blockSize;
    
    __shared__ float smem[blockSize];
    smem[tid] = d_in[gtid];
    __syncthreads();

    for(int index = 1; index < blockDim.x; index *= 2) {
        if ((tid & (2 * index - 1)) == 0){
            smem[tid] += smem[tid + index];
        }
        __syncthreads();
    }
    
    if(tid == 0) d_out[blockIdx.x] = smem[0];
}

这个最直观的实现存在两个严重问题：

1. Warp Divergence（线程分歧）：条件判断导致同一Warp内的线程执行不同路径。随着index增大，活跃线程越来越稀疏，硬件无法高效执行。

2. Bank Conflict（存储体冲突）：当index=16时，线程0访问smem[0]，线程16访问smem[32]，它们属于同一个Bank（Bank 0），导致串行访问。

2.2 V2版本：顺序寻址优化

cpp复制template<int blockSize>
__global__ void reduce_v2(float *d_in, float *d_out){
    __shared__ float smem[blockSize];
    unsigned int tid = threadIdx.x;
    unsigned int gtid = blockIdx.x * blockSize + threadIdx.x;
    
    smem[tid] = d_in[gtid];
    __syncthreads();

    for (unsigned int index = blockDim.x / 2; index > 0; index >>= 1) {
        if (tid < index) {
            smem[tid] += smem[tid + index];
        }