正确性是性能工作的前置条件

GPU bug 往往带有异步、规模相关和非确定性：错误可能在下一个同步 API 才报告；竞态可能只在 release 优化或另一张卡上出现；浮点顺序变化可能被误判为 bug，也可能掩盖真正问题。

建议的证据层次：

CPU/reference implementation
       ↓ 同输入比较
小规模手算 / 单元测试
       ↓
随机、边界、对抗分布
       ↓
Compute Sanitizer（内存/竞态/初始化/同步）
       ↓
多次运行 + 多 block size + 多 GPU 架构
       ↓
端到端不变量与业务指标

Error-checking 基线

所有 Runtime API 都检查返回值；kernel launch 后检查 launch error，并在调试边界同步：

#define CUDA_CHECK(call)                                                     \
  do {                                                                       \
    cudaError_t err__ = (call);                                              \
    if (err__ != cudaSuccess) {                                              \
      fprintf(stderr, "%s:%d CUDA error %d: %s\n", __FILE__, __LINE__,      \
              static_cast<int>(err__), cudaGetErrorString(err__));           \
      std::abort();                                                          \
    }                                                                        \
  } while (0)

#define CUDA_KERNEL_CHECK()                                                  \
  do {                                                                       \
    CUDA_CHECK(cudaGetLastError());                                          \
    CUDA_CHECK(cudaDeviceSynchronize()); /* debug only */                    \
  } while (0)

生产热路径不要在每次 launch 后 device-wide synchronize；应在测试/debug 配置启用，正常路径用 stream/event 建立必要依赖。

异步错误可能导致 context 进入不可继续状态。不要只打印后继续使用可能损坏的数据。

浮点验证

不要逐位比较

并行归约改变加法顺序，FMA 只舍入一次，CPU/GPU 数学库实现和精度路径也可能不同。常用判断：

bool close(float got, float ref, float atol, float rtol) {
  return std::abs(got - ref) <= atol + rtol * std::abs(ref);
}

但容差不是随便给一个 1e-3：

• 根据 dtype、操作次数、数值范围和条件数设置；
• 对接近 0 的结果主要靠 atol，大值主要靠 rtol；
• 单独检查 NaN/Inf、signed zero 和溢出；
• 归约可用 FP64/长双精度 CPU reference 或补偿求和作为对照；
• 必要时比较 ULP 和误差分布，而非只看最大值。

确定性

以下情况常导致运行间差异：

• 浮点 atomic 的竞争顺序；
• 动态任务队列；
• 库根据运行条件选择不同算法；
• 多 GPU collective 的归约树变化；
• 未初始化内存和真实竞态。

前四项可能是合法非确定性，最后一项是 bug。测试合同要明确要求：bitwise、run-to-run、同架构可重现，还是仅数值容差内一致。确定性经常有性能代价。

Compute Sanitizer 四件套

官方 Compute Sanitizer 随 CUDA Toolkit 提供：

推荐编译：

nvcc -O2 -lineinfo app.cu -o app

-lineinfo 让报告映射源码行而不必像 -G 那样关闭许多优化。遇到只在优化版出现的问题，先用接近 release 的构建运行 Sanitizer。

注意边界：

• racecheck 重点是 shared memory hazard，不是所有 global/system 并发 bug 的完整证明器；
• 工具会显著减速，可用 kernel filter、缩小输入和 CI 分层；
• 先跑 memcheck，再处理其他报告，越界可能引发级联假象；
• zero report 不等于程序已被形式化证明正确。

Device assert 与 printf

__global__ void k(const int* index, int n) {
  int i = ...;
  if (i < n) {
    assert(index[i] >= 0);
    if (blockIdx.x == 0 && threadIdx.x < 4)
      printf("i=%d index=%d\n", i, index[i]);
  }
}

仅用于小规模调试：

• printf 使用有限缓冲、扰动时序并严重影响性能；
• 输出通常在同步/launch 边界刷新；
• device assert 失败会使当前工作失败，需要正确处理错误和 context；
• 用 blockIdx/threadIdx 限流，否则日志不可读。

cuda-gdb 与 Nsight Visual Studio Edition

Linux 可用 cuda-gdb 检查 host/device 断点、thread/block、变量和 GPU coredump。Windows 常用 Nsight Visual Studio Edition。优化代码的变量可能被消除或移动；需要源码级单步时用 debug 构建，但最终问题必须回到 release-like 构建验证。

GPU coredump 可由 Compute Sanitizer 或相关环境/API 生成，具体支持和限制按当前官方文档配置。

Nsight Systems 与 Nsight Compute 的分工

Nsight Systems：整机/进程时间线
CPU threads ─ API ─ memcpy ─ kernels ─ synchronization ─ NVTX ─ communication
                         │ 找到热点 kernel / gap
                         ▼
Nsight Compute：单个 kernel 深挖
launch ─ occupancy ─ memory ─ scheduler ─ instruction ─ source/SASS

不要一开始就对整个长程序运行 ncu --set full；计数器回放开销大，且输出会淹没真正热点。

Nsight Systems 工作流

nsys profile \
  --trace=cuda,nvtx,osrt \
  --sample=none \
  --output=report \
  ./app

nsys stats report.nsys-rep

先看：

1. GPU 是否有大段 idle；
2. idle 前 CPU 在做什么：数据准备、锁、同步、launch gap；
3. memcpy 是否与 kernel 重叠；
4. 是否大量短 kernel；
5. 多 stream 是否真的并发；
6. 通信是否成为关键路径；
7. 哪些 kernel 占总 GPU 时间最多。

NVTX 标记业务阶段

#include <nvtx3/nvToolsExt.h>

nvtxRangePushA("preprocess");
preprocess();
nvtxRangePop();

nvtxRangePushA("iteration");
run_iteration();
nvtxRangePop();

没有业务语义的时间线只是一串 mangled kernel 名。用 range 标出 iteration、batch、通信、数据加载和关键算子，可直接回答端到端问题。

Nsight Compute 工作流

先只抓目标 kernel：

ncu --kernel-name regex:my_kernel \
    --launch-skip 10 --launch-count 1 \
    --set full --export report \
    ./app

大型应用可先用较小 section 集合，确认热点后再扩大指标。推荐阅读顺序：

1. GPU Speed of Light：计算与各内存层吞吐接近多少；
2. Launch Statistics：grid/block、寄存器、shared、waves；
3. Occupancy：理论/实测 active warps 及限制因素；
4. Memory Workload Analysis：事务、命中、带宽、bank conflict；
5. Scheduler/Warp State：eligible warps 与 stall；
6. Instruction Statistics：指令混合和管线；
7. Source/SASS：把高成本指标定位到源代码和机器指令。

Profiler 收集某些指标需要 replay kernel，多次执行可能不适合带副作用、依赖随机状态或从外部输入流消费数据的 kernel。查阅当前版本的 replay 模式和限制。

调试与性能构建矩阵

调试时还可设置：

$env:CUDA_LAUNCH_BLOCKING = '1'

它使异步问题更接近发生位置暴露，但改变调度和并发行为，定位完要移除。

测试矩阵

对于同步 bug，改变 block size 和输入尺寸很有效：原本“碰巧同 warp”的线程可能被拆到不同 warp，原本“碰巧同 wave”的 block 也会改变调度。

常见症状速查

官方参考

• Compute Sanitizer
• Nsight Systems User Guide
• Nsight Compute User Guide
• Nsight Compute Profiling Guide
• CUDA Programming Guide: Floating-Point Computation