Reduction 怎么优化？
比如在 GPU 上对一个大数组求和
先弄清楚：每个概念是什么
- Reduction
- 把一大堆数合成一个数，比如 sum、max、dot product。
- Tree reduction
- 两两合并，每一步把剩下的数减半，n 个数 \(\log_2 n\) 步做完。
- Warp shuffle
- 同一个 warp 的 32 个 thread 直接交换 register 里的值，指令是
__shfl_down_sync，不经过 shared memory（见 GPU memory hierarchy）。 - atomicAdd
- 很多 thread 往同一个地址累加时保证一个都不丢，先后顺序不固定。
看图：每一步减半
30 秒回答
I reduce in layers that follow the hardware. First, each thread adds up several elements as it loads them. Then each warp does a tree reduction in registers with shuffle instructions: five steps for 32 lanes, no shared memory, no syncthreads. Next, the warps in a block write partial sums to shared memory and one warp reduces those. Finally the blocks combine with an atomic add or a second kernel. A sum is memory bound, so reading at full bandwidth matters most. Floating-point atomics add in a nondeterministic order, so the last bits can vary.
四层怎么分，手算一遍
对 \(2^{20}\) 个 float 求和，示意的配置是 128 个 block，每个 256 个 thread，一共 32,768 个。
| 层 | 做什么 | 还剩几个数 |
|---|---|---|
| thread | 边读边累加 32 个元素 | 32,768 |
| warp | warp shuffle 做 5 步 | 1,024 |
| block | 8 个 warp 的结果写进 shared memory，一个 warp 再做 3 步 | 128 |
| grid | 128 个 block 用 atomicAdd 累加，或再 launch 一个 kernel | 1 |
warp 这一层的写法。
__device__ float warp_sum(float v) {
for (int off = 16; off > 0; off /= 2)
v += __shfl_down_sync(0xffffffff, v, off);
return v;
}求和每读 4 byte 只做 1 次加法，是 memory bound，所以目标是把 memory bandwidth 用满，相邻 thread 读相邻地址（见 coalescing）。Harris 的讲义在 G80 上分 7 步把 4M 个 int 的求和优化了 30 倍。
必答点
- tree reduction：每步减半，\(\log_2 n\) 步
- 分四层：thread → warp → block → grid
- thread 先边读边累加多个元素
- warp 内用
__shfl_down_sync，不用 shared memory 和__syncthreads - warp 之间经过 shared memory，block 之间用 atomicAdd 或再 launch 一个 kernel
- 浮点 atomicAdd 的顺序不固定，末位会变
追问
shape 是 [M, N]，按行求和。M 很大怎么做，N 很大怎么做？
M 很大时每行分给一个 warp，一个 block 处理很多行。N 很大时一行拆给多个 block，各出一个 partial sum 再合并。
结果和 CPU 上顺序相加的一样吗？
不一定，浮点加法换了顺序末位会不同，验证时按 dtype 设容差（见 kernel 正确性）。要每次结果一致，grid 这一层就再 launch 一个 kernel，代价是多一次 launch。
reduction 前面还有 element-wise 运算怎么办？
融进同一个 kernel，边读边算边累加（见 kernel fusion）。
来源Harris, Optimizing Parallel Reduction in CUDA · NVIDIA blog：Using CUDA Warp-Level Primitives（warp 内求和的写法）· ml-八股 08 第 6 题 · 面经和频率出处见 NVIDIA 题库笔记 1.8 节和 Q2、面试题与 JD 3.2 节