调试与数值正确性：compute-sanitizer 与对齐测试

快但错的 kernel 没有价值。GPU kernel 的 bug 比 CPU 难查——成千上万个线程并发，错误可能只在特定调度下出现、可能静默地给出略错的数值。这一篇讲怎么保证算子正确：用工具查内存和竞态错误，用对参考实现的容差对齐查数值错误，以及怎么理解浮点的不可复现。本系列每个计算 kernel（GEMM、FlashAttention、softmax）都做了这种校验。

一、两类 bug：内存/竞态 与 数值

GPU kernel 的错误大致两类，排查手段不同：

• 内存与竞态错误：越界访问、未初始化读、共享内存数据竞争、漏掉的同步。这类错误可能不报错就给出垃圾结果，或只在某些线程数/输入下崩。靠 compute-sanitizer 这类工具查。
• 数值错误：算法/索引写错导致结果偏差、精度不足、溢出、归一化错误。这类靠和可信参考实现做容差对齐查。

两类都要查——一个 kernel 可能不越界但算错，也可能数值公式对但有竞态。

二、compute-sanitizer：查内存与竞态

compute-sanitizer（CUDA 工具链自带，取代旧的 cuda-memcheck）有几个子工具：

compute-sanitizer --tool memcheck   ./app   # 越界、未对齐、非法地址访问
compute-sanitizer --tool racecheck  ./app   # shared memory 数据竞争
compute-sanitizer --tool initcheck  ./app   # 读未初始化的 global memory
compute-sanitizer --tool synccheck  ./app   # __syncthreads 误用（如发散同步）

• memcheck：抓 第 04 篇 说的越界——边界检查 if (i < n) 漏了、或索引算错读到数组外。它能给出出错的线程和地址。
• racecheck：抓 第 05 篇 的 shared memory 竞争——两个线程不经同步地读写同一 shared 地址。漏 __syncthreads() 是最常见来源。
• synccheck：抓 __syncthreads() 放在只有部分线程进入的分支里（第 04 篇 提过会死锁/未定义）。

开发期应该把它们当 CI 的一部分跑。注意 sanitizer 会显著拖慢执行，所以用小输入跑。另外，WSL2 环境对 compute-sanitizer 的支持并不完整：部分子工具（如 racecheck）可能因底层计数器/驱动支持受限而无法运行或报告不全，需要严格排查竞态时建议在原生 Linux 上复核。

三、对参考实现做容差对齐

数值正确性靠和可信参考比对。标准做法：

1. 用高精度参考：在 CPU 上用 numpy（FP64）算出参考结果，作为”真值”。
2. 容差比较：GPU 用 FP32/FP16 算，不可能和 FP64 逐位相同，要用相对+绝对容差（np.allclose(out, ref, rtol, atol)）而非精确相等。
3. 测多种规模和边界：小尺寸、非 2 的幂、刚好不整除 tile 的尺寸、含 0/负数/大值的输入。

本系列的实践就是这样。例如 GEMM 篇 的四个版本都先在 \(256^3\) 对 A@B 做 allclose 校验通过，再 benchmark。FlashAttention 篇 把 kernel 输出和 numpy 的标准注意力比，最大绝对误差 \(4.12\times10^{-7}\)——在 FP32 精度内完全等价，这个数字证明了在线 softmax 重标定的正确性，而不是”看起来差不多”。

容差怎么定？FP32 累加的相对误差量级约 \(10^{-6}\sim10^{-5}\)（取决于归约长度），FP16 约 \(10^{-3}\sim10^{-2}\)。容差设得比这个量级略松、但远小于”算法错误会产生的偏差”。如果误差远超预期量级，多半不是精度问题而是逻辑 bug。

四、浮点非结合性：不可复现的”错”

一个常见困惑：同一个归约 kernel，多跑几次结果在末位有微小差异，或换并行配置结果变了——这往往不是 bug，是浮点加法不满足结合律：\((a+b)+c \ne a+(b+c)\)（舍入顺序不同）。

GPU 上并行归约的加法顺序取决于线程调度、block 划分。reduction 篇 提到的 atomicAdd 归约尤其明显——原子加的完成顺序不确定，导致 run-to-run 结果在末位浮动。后果：

• 结果对不上要分清原因：是逻辑错（误差大、系统性偏移），还是浮点顺序（误差在 ULP 量级、随机浮动）。后者用容差比较就能通过。
• 需要可复现时：用确定的归约顺序（两趟 kernel 而非原子加）、固定并行配置。训练里为了可复现常牺牲一点性能换确定性。
• 灾难性抵消：相近大数相减丢失有效位。LayerNorm 用 Welford（第 13 篇）而非 \(E[x^2]-E[x]^2\) 正是为避免这个。

五、常见 kernel bug 清单

实战中反复出现的几类：

• 漏边界检查：线程数向上取整后多出的线程越界。症状：memcheck 报越界，或尾部结果错。
• 漏同步：写 shared 后没 __syncthreads() 就读别人写的；racecheck 能抓。
• 发散同步：if (cond) __syncthreads(); 不是全员进入；synccheck 能抓。
• 索引/转置写反：行列搞反、stride 算错。症状：数值系统性错乱，allclose 大幅失败；用小尺寸打印对比定位。
• 累加精度不足：长归约用 FP16 累加、INT8 用 INT16 累加溢出（第 17 篇）。症状：误差随规模增大。
• 未初始化累加器：acc 没清零。症状：结果含垃圾值，initcheck 可辅助。
• 常量/scale 广播错：per-channel scale 维度对错（第 17 篇）。

六、调试方法论

flowchart TD
  A["kernel 结果错"] --> B["先跑 compute-sanitizer<br/>排除越界/竞态/同步"]
  B -->|有报错| C["修内存/同步问题"]
  B -->|干净| D["对 numpy 参考 allclose"]
  D -->|误差巨大/系统性| E["逻辑/索引 bug<br/>小尺寸打印逐元素对比"]
  D -->|误差 ULP 量级浮动| F["浮点非结合性<br/>用容差/确定顺序"]

实用顺序：先用 sanitizer 排除内存和竞态（这类最隐蔽、最该先排），再用容差对齐查数值，最后区分逻辑错与浮点顺序。小尺寸 + printf（kernel 内可用）打印中间值，是定位索引类 bug 最直接的手段。

七、小结与下一步

• GPU bug 分内存/竞态和数值两类，分别用 compute-sanitizer 和容差对齐排查。
• compute-sanitizer 的 memcheck/racecheck/synccheck/initcheck 分别查越界、shared 竞争、同步误用、未初始化读，应纳入 CI。
• 数值正确性对 FP64 numpy 参考做 allclose；本系列 FlashAttention 实测误差 4e-7 即等价证明。容差按精度量级设。
• 浮点不满足结合律，并行归约的末位浮动通常不是 bug；需可复现时用确定顺序。

正确性是底线。最后一篇展望未来——趋势：TMA、Blackwell、ThunderKittens 与编译器协同，看算子工程往哪走。