CUDA 归约求和（Reduction）算法
这段代码是CUDA并行编程中经典的归约求和（Reduction）算法的核心部分，用于高效累加数组元素。下面从零开始逐步解析：
1. 基础概念
- 目标：将数组所有元素求和（例如
[3,1,4,2]→10）。 - 并行策略：通过多线程分层累加，逐步缩小计算规模。
2. 代码逐行解析
**(1) 初始化步长**
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
blockDim.x：当前线程块中的线程总数（如256）。- 初始化为线程数的一半（如128），每次循环右移1位（即除以2）：
- 迭代顺序：128 → 64 → 32 → … → 1。
**(2) 条件判断**
if (tid < stride)
tid：当前线程的ID（0 ≤ tid < blockDim.x）。- 仅让前一半线程参与计算（如
tid < 128时）。
**(3) 累加操作**
sdata[tid] += sdata[tid + stride];
- 动作：将后半段元素加到前半段对应位置：
tid=0：sdata[0] += sdata[128]tid=1：sdata[1] += sdata[129]- …
- 效果：数组规模减半（如256 → 128 → 64 → … → 1）。
**(4) 同步屏障**
__syncthreads();
- 确保所有线程完成当前步长的累加，再进入下一轮。
**3. 示例演算（4元素归约）**
假设 sdata = [3,1,4,2]，blockDim.x=4：
- 第1轮（stride=2）：
tid=0：sdata[0] += sdata[2]→3+4=7tid=1：sdata[1] += sdata[3]→1+2=3- 结果：
[7,3,4,2]（后两元素不再使用）
- 第2轮（stride=1）：
tid=0：sdata[0] += sdata[1]→7+3=10- 结果：
[10,3,4,2]（最终和=10）
4. 关键设计思想
- 分层合并：每次迭代将数据规模减半，时间复杂度从O(n)降至O(log n)。
- 线程复用：仅活跃线程数逐轮减少，避免资源浪费。
- 共享内存：
sdata需用__shared__声明，实现线程块内高速通信。
5. 完整CUDA归约示例
cppCopy Code__global__ void reduceSum(int* input, int* output) {extern __shared__ int sdata[];int tid = threadIdx.x;sdata[tid] = input[tid];__syncthreads();// 归约核心代码for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {if (tid < stride) {sdata[tid] += sdata[tid + stride];}__syncthreads();}if (tid == 0) output[0] = sdata[0]; // 写入结果
}
6. 常见问题
- **为什么需要
__syncthreads()？**
防止线程间数据竞争（如某线程未完成写入，其他线程已读取）。 - **如果数组不是2的幂次？**
需填充虚拟元素或特殊处理边界条件。
通过这种分治策略，GPU可高效处理大规模数据归约，广泛应用于统计、机器学习等场景。
1. 执行范围
- 所有线程均会执行：线程块（Block）内的所有线程都会进入这段代码（从
threadIdx.x=0到threadIdx.x=blockDim.x-1）911。 - 条件分支：只有满足
threadIdx.x < stride的线程会执行加法操作，其他线程跳过加法但**仍会执行__syncthreads()**。
2. 关键行为分析
**(1) 条件执行部分**
cppCopy Codeif (threadIdx.x < stride) {sdata[threadIdx.x] += sdata[threadIdx.x + stride];
}
- 活跃线程：仅
threadIdx.x < stride的线程执行加法（如stride=256时，前256个线程活跃）。 - 非活跃线程：虽不执行加法，但仍需参与后续同步。
**(2) 同步部分**
__syncthreads(); // 所有线程必须到达此点才能继续
- 强制同步：线程块内所有线程必须执行到此语句并等待其他线程，确保共享内存