7. 性能测量、数值与常见陷阱

GPU 上的性能直觉比 CPU 更容易错：同一段代码，改一个 block 大小可能差 3 倍，换个访问顺序可能差 8 倍。这一节先把“怎么测”定下来，再给三个可以手算的极限，最后是移植检查清单。

7.1 怎么测才算数

三步，缺一不可：warmup、多次采样、每次之间加同步。

bench.cu

#include <algorithm>
#include <vector>

// 1) warmup：让时钟爬升、cache 预热、驱动完成首次加载
for (int i = 0; i < 10; ++i) kernel<<<blocks, 256>>>(d, n);
CK(cudaDeviceSynchronize());

// 2) 逐次计时，拿分布而不是一个数
cudaEvent_t a, b;
CK(cudaEventCreate(&a)); CK(cudaEventCreate(&b));
std::vector<float> ms(iters);

for (int i = 0; i < iters; ++i) {
    CK(cudaEventRecord(a));
    kernel<<<blocks, 256>>>(d, n);
    CK(cudaEventRecord(b));
    CK(cudaEventSynchronize(b));          // 只等这一次
    CK(cudaEventElapsedTime(&ms[i], a, b));
}

std::sort(ms.begin(), ms.end());
printf("median %.4f ms, p90 %.4f ms\n", ms[iters / 2], ms[iters * 9 / 10]);

为什么要这样：

• 第一次运行不算数。GPU 会从低功耗状态爬频，驱动也可能在第一次 launch 时做 JIT 和模块加载。
• 中位数比平均值可靠。偶发的上下文切换和时钟波动会把平均值拉偏；p90 能告诉你尾延迟。
• 同步的位置要对。测吞吐时在循环外同步一次；测单次延迟时每次同步。两者不能混。
• 用 debug 编译测性能没有意义。-G 会关掉几乎所有优化，数字可以差 10 倍。

还要留意降频：nvidia-smi -q -d CLOCK 可以看到当前频率是否掉到 base clock。功耗墙和温度墙在长时间跑大 kernel 时很常见。

7.2 三个极限：先算下界，再看差距

任何 kernel 的时间都不可能低于这三者中的最大项：

1. 带宽下界：必须搬运的字节数除以 \(BW\)。
2. 算力下界：必须做的浮点运算数除以 \(P\)。
3. 延迟下界：启动开销、同步点、依赖链长度之和。

先手算前两个，取较大者，再和实测比。如果实测是下界的 1.1–1.5 倍，已经不错；如果是 5 倍，先查访问模式和占用率。下面这个例子把三个都算一遍：

最后一行解释了为什么“把 100 次 launch 合成 1 次、在 kernel 里循环”经常能立刻提速几倍。启动开销是固定成本，和 kernel 大小无关。

有效带宽的定义是

例如 1 GiB 的拷贝（读 + 写共 2 GiB）耗时 1.2 ms，则 \(BW_{\text{eff}}\approx1.7\) TB/s，约为 H100 峰值的 50%。低于 50% 时，先看合并访问（2.3）和每指令字节数（2.2）。

7.3 profiler 看哪几个数

不要靠猜。Nsight Compute（ncu）先看“Speed of Light”页的两个百分比：

• Memory Throughput：显存带宽用了多少。接近 100% 说明带宽受限，优化方向是减少流量。
• Compute (SM) Throughput：算力用了多少。接近 100% 说明算力受限，优化方向是指令效率（减少冗余计算、用 FMA、用 Tensor Core）。
• Achieved Occupancy：实际常驻 warp / 上限，和 1.5 的估算对照，能发现寄存器或 shared 超预期。
• Warp Stall Reasons：long scoreboard 高 = 等显存；barrier 高 = 同步太多或负载不均；MIO throttle 高 = shared/特殊函数挤爆了流水线。
• L2 Hit Rate 与 Sectors/Request：后者接近 4（一个 warp 128 B = 4 个 sector）说明合并良好；明显大于 4 就是访问散了。

Nsight Systems（nsys）看时间线：传输和计算有没有重叠、流之间有没有空隙、有没有意料之外的同步。它回答“整体时间花在哪”，Compute 回答“这个 kernel 内部为什么慢”。

正确性用 compute-sanitizer：--tool memcheck 查越界，--tool racecheck 查 shared 竞争，--tool initcheck 查未初始化读。它的开销很大，只在调试时开。

7.4 数值：精度、FTZ 与确定性

和 CPU 相同的是：都遵循 IEEE 754 的舍入规则，误差用 ULP 或相对误差衡量，Kahan 补偿求和同样有效。不同的是：

• FTZ：GPU 默认可能把非规格化数刷成 0（-ftz=true），CPU 一般保留。极小值附近的差异不是 bug。
• --use_fast_math 会把除法、开方、expf 换成近似指令，误差可达几个 ULP 到 \(10^{-4}\)，还会打开 FTZ。用它之前先确认容差。
• 累加顺序：归约树、原子、Tensor Core 的累加顺序都和串行不同，结果无法逐位复现。跨实现对比用相对误差（CPU 侧的门槛通常是 \(10^{-6}\)，Tensor Core 用 \(10^{-3}\)）；要固定顺序的写法见 4.1。
• 原子浮点求和：顺序由调度决定，每次运行都可能不同。要复现就固定树形结构，或改用整数/定点。

7.5 移植检查清单

1. 数据量够大吗？低于几十万元素、或只跑一次，传输和启动就吃掉收益（0.5）。
2. 数据能常驻显存吗？能，才谈得上加速；每次请求都重传，先算 PCIe 的账。
3. 分支整齐吗？同一 warp 内路径是否一致；能不能改写成无分支（查表、掩码、排序）。
4. 访问合并吗？看每个 warp 覆盖的字节数、Sectors/Request、AoS 能否改 SoA。
5. 复用进 shared 了吗？tile 是否够大（≥ 64×64），有没有 bank conflict。
6. 占用率与 ILP 够吗？寄存器有没有爆、每线程是否只发一个 4 B 请求（0.3）。
7. 启动与传输占比多少？小 kernel 合并，流式传输用 pinned + 多流（6.2、6.3）。
8. 精度在容差内吗？换了 Tensor Core 或 fast math 之后重新验收。
9. 需要复现吗？去掉原子汇总，固定归约顺序。
10. 用 profiler 验证过瓶颈吗？没有数据支撑的优化都是猜。

7.6 容易写错的地方