浮点数精度问题基础
做异构计算一致性检查，绕不开一个老熟人——浮点数精度。我第一次遇到GPU和CPU结果对不上时，第一反应是代码写错了。排查了一整天，最后发现是两个平台对同一个浮点表达式的计算路径不同。说白了，这个问题从IEEE 754标准诞生那天就埋下了。
2.1 IEEE 754 标准速览
IEEE 754定义了单精度（float32）和双精度（float64）的二进制表示。我记得刚入行时对着那张格式图看了好久才明白，符号位、指数、尾数都是在有限比特里做权衡。
| 类型 | 总位数 | 符号位 | 指数位 | 尾数位 | 约十进制有效数字 |
|---|---|---|---|---|---|
| float16（半精度） | 16 | 1 | 5 | 10 | 3~4位 |
| float32（单精度） | 32 | 1 | 8 | 23 | 7~8位 |
| float64（双精度） | 64 | 1 | 11 | 52 | 15~16位 |
| bfloat16（Google脑格式） | 16 | 1 | 8 | 7 | 2~3位 |
为什么会有多种格式？讲白了，性能取舍。GPU用float32跑得快，但做累加时误差累积比double严重。我在项目里见过一个分子动力学模拟，跑了100万步后GPU结果和CPU差了好几个分子——根源就是累积误差。
2.2 精度损失的典型场景
举个例子：
float a = 1.0000001f;
float b = 1.0000000f;
float c = a - b; // 理论值 0.0000001
// 实际上 c 可能变成 1.0e-7 左右，但精度只有7位有效数字
在GPU上做归约求和时，这种情况尤其常见。每个线程局部求和后，主线程再做一次加法。如果中间结果差异很大，精度就崩了。
我的避坑经验：我曾经写过一个向量点积kernel，CPU算出来是0.00123456，GPU算出来是0.00123478。差在最后两位。排查发现是加法顺序不同导致舍入不同。解决方法很简单——改用Kahan求和算法（一种补偿求和法），两边结果就对齐了。
2.3 运算顺序与结合律
实数上的结合律 (a+b)+c = a+(b+c)，在浮点数世界里不成立。因为每一步都会舍入。GPU的并行归约本质上是树形加法，与CPU顺序加法的顺序完全不同。
// 顺序求和（CPU风格）
float sum_seq = 0;
for (int i=0; i<N; ++i) sum_seq += data[i];
// 树形归约（GPU风格）
// 每步两两相加，结果写回共享内存
// 最终只有一条路径到根
两种运算路径不同，舍入误差也不同。我在调试一个天气模型时，发现GPU版和CPU版在某个经纬度网格上的温度差了0.02K。最后用-ffast-math编译选项把CPU的运算顺序改成了类似GPU的树形归约，误差立马降到了0.0001K以内。
-fmad=true做融合乘加（FMA），使运算路径更接近CPU。
2.4 FMA（融合乘加）的陷阱
FMA执行 a*b + c 时只舍入一次，而不是逐次舍入。这对精度是好事，但对一致性检查是坏事——GPU和CPU可能一个用FMA一个不用。
我在一个推荐系统里遇到过：矩阵乘法的结果在GPU和CPU上差了1e-6。排查后发现，GPU的CUDA核心默认开启FMA，而CPU的编译器没开。解决方案是：编译CPU代码时加上 -mfma（x86）或 -ffp-contract=fast（gcc），让两边行为一致。
# 编译CPU代码时启用FMA
gcc -O3 -mfma -ffp-contract=fast main.c -o main
# 或者反过来，禁用GPU的FMA
// CUDA代码中：
#pragma nv_exec_check_disable
__global__ void my_kernel() {
// 使用 __fmul_rn 和 __fadd_rn 替代自动FMA
}
2.5 不同精度混合计算
GPU上经常混用float16、float32、float64。比如矩阵乘法用float16做核心计算，用float32累积。CPU上可能全程用double。两边的中间表示不同，结果自然不同。
| 混合模式 | GPU常见做法 | CPU常见做法 | 一致性风险 |
|---|---|---|---|
| 输入fp16，累加fp32 | Tensor Core自动转换 | 手动cast后累加 | 中 |
| 输入fp32，中间fp32 | 默认模式 | 默认模式 | 低（顺序差异仍存在） |
| 输入fp32，中间fp64 | 需显式声明double | 默认double更容易 | 高（精度差太多） |
个人习惯：我在写一致性测试用例时，会先统一两边的精度配置。比如都强制用float32做计算和累积，不许编译器自动提升到double。这样至少能排除精度差异带来的干扰。