I·Infra Notes AI Infra 工程实践指南 ← 返回主博客
GPU 与算子 / READING NOTES

2.2 Reduction、Scan 与 Softmax 的共同骨架

CUDA 索引、访存、归约、扫描与 Softmax

每线程处理若干元素 → Block 内树形合并 → 跨 Block 二次归约
Scan 多一个“前缀传播”阶段;Softmax 用 max 与 sum 两次归约
import torch

def stable_softmax(x, dim=-1):
    m = x.max(dim=dim, keepdim=True).values
    e = torch.exp((x - m).float())
    return e / e.sum(dim=dim, keepdim=True)

def inclusive_scan(x):
    return torch.cumsum(x, dim=-1)

CUDA Block 归约伪代码(每步同步;最后一个 Warp 可用 shuffle):

shared[threadIdx.x] = local_sum;
__syncthreads();
for (int offset = blockDim.x / 2; offset > 0; offset /= 2) {
    if (threadIdx.x < offset) shared[threadIdx.x] += shared[threadIdx.x + offset];
    __syncthreads();
}
if (threadIdx.x == 0) partial[blockIdx.x] = shared[0];

优化路径:先测试非 2 的幂长度、空行、极大值和不同 dtype;再减少跨 Block 原子争用、把多元素归约放到寄存器,检查寄存器压力与同步次数。Scan 不能把各 Block 独立扫描后直接拼接,必须把前序 Block 总和加到后序输出。

Bug 案例:只在部分线程进入 __syncthreads() 可能死锁;Softmax 忘记减 max 会溢出为 NaN;未 mask 的 padding 会改变分母。追问:在线 Softmax 需要保存哪两个统计量?历史最大值 m 与指数和 l。