跳到主要内容
TileLang 系统教程

tilelang-tutorial.huandu.me

第 12 章 调试与性能分析

/code/12-debugging-and-profiling/

02-tolerance-model/README.md2822 字节
# 示例 2:低精度对拍的容差从哪来

## 目的

正确性对拍最后一定会落到一行断言:`atol`、`rtol` 取多少?「±1e-3」这种数字如果是
抄来的,出了问题既不敢放宽也不知道能收紧。这个脚本把容差拆成三层:

1. **机器精度**:从 IEEE-754 的有效位数算出 fp16 / bf16 / fp32 的 `u = 2^-p`;
2. **归约预算**:长度 K 的 dot / GEMM-K 维,输入舍入误差相对 `S = Σ|a_i·b_i|`
   的最坏情形是 `u`,均方根是 `u/√K`;
3. **实测分布**:`fp16/bf16 输入 + fp32 累加` 的 dot 与 fp64 精确值比较,看实测
   中位数与预算是否同量级。

它证明的是「容差的数量级从哪来」,**不是**「某个 kernel 一定通过某个容差」——后者
取决于你的数据分布、归约顺序和累加类型,必须在有卡机器上实测。

## 运行前提

* Python 3.9+ 与 numpy(TileLang 自身也依赖 numpy)。
* 不需要 GPU/NPU。

## 运行命令

```bash
python3 chapters/12-debugging-and-profiling/code/02-tolerance-model/tolerance_model.py
python3 tolerance_model.py --lengths 128,512,2048 --trials 4000
```

## 期望输出

作者在写作机上的实际输出(2026-10-01,Python 3.11 + numpy 1.24.1,默认 2000 次试验):

```
== 一、由有效位数得到的机器精度 ==
格式    有效位 p     u = 2^-p      10 的幂提示 atol    说明
fp16    11          4.883e-04       1e-03             binary16:1 符号 + 5 指数 + 10 尾数
bf16    8           3.906e-03       1e-02             bfloat16:1 符号 + 8 指数 + 7 尾数
fp32    24          5.960e-08       1e-07             binary32:1 符号 + 8 指数 + 23 尾数

== 三、实测:低精度输入 + fp32 累加 vs fp64 精确值 ==
每个 K 跑 2000 次随机试验(固定种子 20261001)。
格式    K       实测最大            实测 99.9 分位        实测中位数           RMS 预算        中位数/预算
fp16    64      1.521e-04         1.357e-04           2.984e-05         6.104e-05       0.49
fp16    256     7.305e-05         6.968e-05           1.481e-05         3.052e-05       0.49
fp16    1024    3.514e-05         3.320e-05           7.764e-06         1.526e-05       0.51
fp16    4096    2.336e-05         1.909e-05           3.704e-06         7.629e-06       0.49
bf16    64      1.383e-03         1.095e-03           2.366e-04         4.883e-04       0.48
bf16    256     5.791e-04         5.164e-04           1.135e-04         2.441e-04       0.46
bf16    1024    3.100e-04         2.927e-04           5.901e-05         1.221e-04       0.48
bf16    4096    1.688e-04         1.478e-04           2.906e-05         6.104e-05       0.48
```

**判据**:「中位数/预算」这一列稳定在 0.5 左右,说明 `u/√K` 这个量级模型可用;
若它偏离一个数量级以上,说明数据分布或者累加类型与模型假设不符。

## 这张表怎么用在真实对拍里

* 上表的基准是 `S = Σ|a_i·b_i|`。日常断言用的是「相对结果」口径,
  当数据零均值时 `S / |Σa_i·b_i| ≈ √K`,所以长归约的相对误差可以放大到 `u·√K`
  (fp16、K=4096 约 3e-2),**因此长归约必须配 atol 兜住接近零的输出**。
* 工程起点:fp16 用 `rtol=1e-3, atol=1e-3`;bf16 用 `rtol=1e-2, atol=1e-2`;
  再配 `max_mismatched_ratio≈1e-3` 允许极少数离群元素。
* 这与锁定版本里两个对拍工具的缺省量级一致:
  `torch_assert_close(rtol=1e-2, atol=1e-3, max_mismatched_ratio=0.001)` 与
  `Profiler.assert_allclose(atol=1e-2, rtol=1e-2, max_mismatched_ratio=0.01)`。

## 注意事项

* 本脚本用「截掉 fp32 低 16 位尾数」近似 bf16 舍入,与硬件的 round-to-nearest-even
  有微小差异,只影响第三位有效数字,不影响数量级结论。
* 这只覆盖了「输入舍入」这一主项。真实 kernel 还可能有累加顺序、分块累加(split-K)、
  以及 epilogue 中额外的 cast 误差,需要在有卡机器上用固定输入实测。

## 作者实际做到的验证级别

**CPU 逻辑验证**:在 macOS (Apple Silicon)、Python 3.11 + numpy 1.24.1 上实际运行过,
退出码 0。未在任何加速器上运行过任何 kernel。

← 浏览本章全部代码