浮点数与数值稳定性：从 IEEE 754 到 FP8 的 LLM 精度世界

从十进制科学计数法说起#

你早就熟悉十进制科学计数法：$12345.6 = 1.23456 \times 10^4$。它把一个数拆成两部分：有效数字（1.23456）和指数（4），有效数字固定在 $[1, 10)$ 区间内，叫做”规格化”（normalized）形式。

二进制完全同理：任何非零实数 $x$ 都可以唯一写成

其中 $f$ 是小数部分（二进制），$e$ 是整数指数。比如 $5.5_{10} = 101.1_2 = 1.011_2 \times 2^2$。关键差异只有一个：指数部分是 $2$ 的幂而不是 $10$ 的幂。

那么存储一个数需要哪些信息？三个：符号（正负）、指数 $e$、小数 $f$。IEEE 754 标准（1985 年首次发布，现行版本 IEEE 754-2019）就用固定位数把这三个部分装进一个字节串。以单精度 FP32 为例，它占 32 位：

• 最高位 1 位：符号位 $S$，$S=1$ 表示负数；
• 中间 8 位：指数位，存放 $e$ 加一个偏移量（bias）；
• 低 23 位：尾数位，存放规格化形式里小数部分 $f$ 的二进制位。

于是数值为

其中 $E$ 是指数位组成的无符号整数。IEEE 754 对位布局的每一个细节都值得单独解释，因为每个细节都对应一个数值特性。

为什么指数要加偏移量（bias）而不是用补码#

FP32 的 8 位指数可以表示 $0 \sim 255$ 的无符号整数。如果我们直接把 $e$ 存成带符号数，比较两个浮点数大小时还得先处理符号位和指数的符号，硬件设计变复杂。IEEE 754 的做法是存 $E = e + 127$，即偏移量为 127。这样指数部分 $E$ 在内存里就是”越大代表数越大”的无符号数，配合符号位和尾数，整个浮点数可以用无符号整数比较来完成大小比较（这叫 monotonic ordering，是 float 排序和 radix sort 的基础）。这也意味着最小的正规格化数是 $e = -126$（$E = 1$），最大的指数是 $e = 127$（$E = 254$）。

为什么尾数不用存开头的那个 1#

规格化形式里 $x = 1.f \times 2^e$ 的整数部分永远是 1（对正规格化数而言），那这一位就不必存储——这叫”隐式前导位”（implicit leading bit）。23 位尾数字段因此白赚了 1 位精度：FP32 实际有效数字是 24 位。同理 FP16 有 11 位有效数字、BF16 只有 8 位。

特殊值：指数和尾数全 0 / 全 1 时#

指数位全 0 或全 1 被 IEEE 754 保留给特殊值：

次正规数值得单独说明：最小规格化 FP32 数是 $2^{-126} \approx 1.18 \times 10^{-38}$。如果比它还小的正数全部变成 0，那么在 $0$ 和 $2^{-126}$ 之间就出现”断层”，减法 a - b 在 $a \approx b \approx 10^{-40}$ 时可能直接得 0，破坏”减法结果平滑过渡”的性质。次正规数用”尾数前导位为 0”的方式把最小可表示数扩展到 $2^{-149} \approx 1.4 \times 10^{-45}$，代价是这些数的精度随指数下降而下降。GPU 上次正规数运算通常有额外开销（CUDA 里可以用 --ftz=true 把次正规数 flush 成 0 换性能，代价是精度）。

浮点数在数轴上是”疏密不均”的#

一个重要的直觉：浮点数不是均匀分布在数轴上，而是越靠近 0 越密集，越往两边越稀疏。在 $[1, 2)$ 区间内，FP32 的间距是 $2^{-23} \approx 1.19 \times 10^{-7}$；在 $[2, 4)$ 间距翻倍成 $2^{-22}$；到 $[2^{20}, 2^{21})$ 间距已经是 $2^{-3} = 0.125$。也就是说，浮点数的相对精度（有效数字位数）是恒定的，绝对精度随数值大小线性增长。这是浮点表示与定点/整数表示的本质区别，也是后文一切讨论的出发点：浮点数的误差是”相对误差”，不是”绝对误差”。

表示误差：0.1 在二进制里是无理数#

0.1 + 0.2 ≠ 0.3 的原因现在就清楚了：$0.1_{10}$ 的二进制展开是无限循环小数 $0.0001100110011\cdots_2$，任何有限位数的浮点格式都无法精确表示它，只能存最接近的一个值。这个”表示误差”（representation error）约等于该数附近的间距的一半。0.1 与 0.2 各自都有约 $10^{-17}$ 量级的表示误差，相加后误差没被抵消，于是结果和”0.3 的最接近表示”差了一个最小单位（ULP，unit in the last place），打印出来就是 0.30000000000000004。

这种误差本身无害（相对误差只有 $10^{-16}$ 量级），但它揭示了浮点运算的第一定律：每个浮点数都只是真实值的近似，每次运算都可能引入新的近似误差。数值分析（numerical analysis）这门学科研究的正是这些误差如何传播、何时会毁掉结果。

机器精度 $\epsilon$：1 和下一个数的距离#

最常用的误差尺度叫机器精度（machine epsilon）$\epsilon$：数值 1 与下一个可表示浮点数之间的间距。

（严格说这是”spacing at 1”，有的教材把”半间距”也叫 epsilon；本文统一用间距。）$\epsilon$ 告诉你”相对精度大概是 $10^{-7}$“意味着你的数有约 7 位十进制有效数字。FP16 有约 3 位，BF16 约 2 位，FP8 的 E4M3 只有 1 位左右。每次浮点运算的相对误差上界大约是 $\epsilon/2$（半间距，round-to-nearest 下）。

舍入模式：默认是”最近偶数”#

浮点运算结果往往不是可表示的数，必须舍入。IEEE 754 默认模式是 round-to-nearest, ties-to-even：舍入到最近的可表示数，恰好在中点（tie）时选尾数为偶数的那个。选”偶数”而不是”向上”是为了避免统计偏差——如果总是向上舍入，一万次运算的误差会系统性偏正。随机误差可以互相抵消，系统性偏差则会累积成可观的错误。

大数加小数：1 + 2^{-24} = 1 的完整推导#

现在来精确分析那个经典的例子。在 FP32 里：

• $1$ 的可表示形式是 $1.000\cdots0_2 \times 2^0$，下一个数是 $1.000\cdots1_2 \times 2^0$（尾数最低位加 1），间距 $\epsilon = 2^{-23}$。
• 要加上的 $2^{-24}$ 是多少？它是 $2^{-24}$，恰好等于半间距 $\epsilon/2 = 2^{-24}$。
• 结果 $1 + 2^{-24}$ 正好落在 $1$ 和 $1 + 2^{-23}$ 的正中间——tie。
• ties-to-even 规则：$1.0$ 的尾数全 0（偶数），$1 + 2^{-23}$ 的尾数最低位是 1（奇数），所以舍入到 $1.0$。

结论：FP32 中 1.0f + 1e-8f（约 $2^{-26.6}$，比半间距还小）必然等于 1.0f。更一般的规律：当两个数 $a > b$ 的指数之差超过有效数字位数时，$a + b$ 直接舍入为 $a$。FP32 有 24 位有效数字，所以比值超过 $2^{24} \approx 1.7 \times 10^7$ 的小数会被完全吞掉。FP16 的 11 位有效数字意味着比值超过 $2^{11} = 2048$ 就吞掉，BF16 是 $2^8 = 256$。

这解释了一个常见的困惑：(1e8 + 1) - 1e8 = 0。在 $10^8 \approx 2^{26.6}$ 附近，FP32 的间距是 $2^{26.6 - 23} \approx 2^{3.6} \approx 12$，1 远小于半间距 6，所以 1e8 + 1 先被舍入回 1e8，减法后自然得 0。注意这里每一步都”正确”地舍入了，错误来自信息在那一步已经被丢弃——这就是数值误差的本质：信息一旦被舍入抹掉，后面任何操作都找不回来。

消去误差（catastrophic cancellation）：相近数相减#

比”大数加小数”更隐蔽的是相近数相减。假设 $a = 1.00000019$，$b = 1.00000004$，精确差是 $1.5 \times 10^{-7}$。但 FP32 里 $a$ 和 $b$ 各自的表示误差约 $5 \times 10^{-8}$（半间距），$a - b$ 的误差上界是两者误差之和 $10^{-7}$——和结果本身 $1.5 \times 10^{-7}$ 同一个量级，相对误差可能高达 50% 以上。

这不是舍入出了错，而是”减法把两个数共有的高位全部抵消，只剩误差”：$a$ 和 $b$ 都有 7 位有效数字，但它们相同的部分（1.000000）在差里消失了，差的相对精度取决于 $a - b$ 的绝对误差，而绝对误差由 $a$、$b$ 的舍入误差决定。公式化地说：

分子是”输入误差”，分母是”结果大小”。当 $a \approx b$ 时分母趋近 0，相对误差爆炸。数值分析里给这个放大倍数一个名字叫条件数（condition number），后面还会遇到。

在 LLM 里消去误差最常见的藏身处是求和的先后：注意力分数、梯度、统计量全都涉及大量求和，而浮点加法不满足结合律——这是下一节的内容。

浮点加法不结合：(a+b)+c ≠ a+(b+c)#

取 $a = 10^8$，$b = -10^8$，$c = 1$：

两条路径结果不同！原因就是 -1e8 + 1 在 FP32 里等于 -1e8（前面分析过，间距 12 吞掉了 1）。浮点加法是可交换的，但不是可结合的——结合律失效直接导致”求和顺序影响结果”，这也是 GPU 上归约（reduction）为什么必须小心设计的原因。

假设要计算 $S = \sum_{i=1}^{n} x_i$。用朴素顺序求和，误差上界大约为：

如果所有 $x_i$ 同号，这个上界基本是紧的——误差随 $n$ 线性增长。用分层求和（tree reduction，两两相加再两两相加，GPU 硬件上正是这么做的）可以把上界降到 $\epsilon \log_2 n \cdot \sum|x_i|$，从线性变成对数。如果 $x_i$ 有正有负，误差呈随机游走，约按 $\sqrt{n} \cdot \epsilon$ 增长，比最坏情况好得多——这也是为什么”不要排序后把所有正数先加完再加负数”。

更进一步的技巧是 Kahan 求和（1965）：用一个补偿变量记录每一步丢掉的低阶误差，下次加法时补回去，误差上界降到 $\sim 2\epsilon \sum|x_i|$，与 $n$ 无关。CUDA 的 cub::WarpReduce、PyTorch 的 torch.sum 内部都实现了分层求和，大部分情况下你不需要手写 Kahan，但要理解”顺序求和精度差、分层求和精度好”这个事实——比如你手写注意力 kernel 时，把 softmax 分母的求和按 block 累加，就要意识到误差随块数增长，最终项全部同号，上界是线性的，长序列下必须用 FP32 甚至 Kahan 来保精度（FlashAttention 的官方实现里，在线 softmax 的求和恰好就是这种情况）。

条件数：从”误差放大”到线性代数#

把误差放大倍率形式化就得到条件数。对线性方程组 $Ax = b$，解的相对误差满足：

其中条件数 $\kappa(A)$ 衡量”输入扰动被放大的倍数”。它与 SVD（奇异值分解）直接相关——SVD 的分解式、最优截断定理（Eckart–Young）等完整讲解见《答疑特辑：秩与 SVD、RoPE 旋转矩阵、投机解码原理与 AI 芯片全景》：

即最大奇异值与最小奇异值之比——这正好回答”奇异值大小对应矩阵的什么特性”：奇异值刻画矩阵在每个方向上的拉伸倍数，最大与最小的比值就是这个矩阵”有多接近退化”的度量。$\kappa$ 大意味着矩阵把某个方向拉伸得极长、另一个方向压得极扁（几何上接近”退化”），这样的矩阵求逆时，微小的输入误差会被放大 $\kappa$ 倍。$\kappa = 10^{14}$ 的矩阵在 FP32（$\epsilon \approx 10^{-7}$）下求逆，解的相对误差理论上可达 $10^7$——结果基本是噪声。实际数值算法里通常用 cond() 预估误差，或者干脆避免显式求逆。这个”输入误差 × 条件数 = 输出误差”的框架，是理解数值稳定性问题的统一语言：一个算法数值稳定，指的就是它的输出对输入的相对误差放大倍数（有效条件数）不失控。

一张表看全所有格式#

（E4M3 的 448 这个”不整齐”的数字来自它砍掉 Infinity 的做法，后面细讲；表格里”有效数字”包含隐式前导位。）

BF16 为什么只有 8 位有效数字却统治了训练#

BF16 的位布局是 $1 + 8 + 7$：指数位数与 FP32 完全相同（8 位），只把尾数砍到 7 位。它由 Google 在 2019 年的论文 A Study of Bfloat16 for Deep Learning Training 中系统提出（TPU 早在 2016 年就开始用了）。论文的核心实验结论是：在 DeepSpeech2（语音）、GNMT（翻译）、SRGAN（图像）等覆盖 CNN/RNN/Transformer 的工作负载上，BF16 训练与 FP32 训练收敛到几乎相同的精度，而内存占用减半。

为什么砍尾数没事、砍指数会出事？因为深度学习数值的”危险模式”是溢出/下溢，不是精度不足：中间值（尤其是梯度）的动态范围跨越几十个数量级，FP16 的最大值只有 65504，梯度一放大就溢出成 Infinity，训练直接炸掉；而 BF16 的范围和 FP32 一模一样，永远不会因为范围问题翻车。牺牲的只是 2-3 位十进制精度，而深度学习对权重/激活的精度需求恰好没那么苛刻（大多数情况下 $10^{-2}$ 量级的相对误差都扛得住）。这就是”范围 > 精度”的取舍逻辑，它后来成了低精度计算的第一原则。

FP16 的悲剧：更”精密”却更难用#

FP16 有 5 位指数、10 位尾数，有效数字比 BF16 多 3 位，但动态范围只有 $6.1 \times 10^{-5} \sim 65504$——上界比 FP32 小 33 个数量级，下界又不够小。训练时 FP16 要么溢出（大梯度）要么下溢（小梯度），所以必须配合 loss scaling 技巧（把 loss 乘一个大常数，把梯度”抬”进可表示范围）。混合精度训练（mixed precision，FP16 计算 + FP32 主权重 + loss scaling）比 BF16 方案复杂得多，这也是后来 NVIDIA 在 Hopper 之后主推 BF16、FP16 逐渐退居存储格式的原因。推理侧 FP16 相对更常见（权重幅度已知，不会突然爆炸），但激活仍然有风险——RMSNorm 之后的值如果掉到 $6 \times 10^{-5}$ 以下就直接变成 0，某些小方差层会”莫名其妙丢精度”。

TF32：A100 张量核心上的”截尾 FP32”#

TF32 是 NVIDIA 随 Ampere A100 引入的格式：指数 8 位（和 FP32/BF16 相同）、尾数只有 10 位（和 FP16 相同），总共 19 位。它本质上是把 FP32 的 23 位尾数截断成 10 位的”精度打折版 FP32”，专门在张量核心上跑矩阵乘。好处是硬件上直接复用 FP32 的指数通路，范围与 FP32 完全一致，输入输出还是 FP32（无需额外转换），坏处是每次乘法的相对误差放大到 $2^{-11}$ 左右，相当于 FP16 的精度。A100 上 TF32 矩阵乘吞吐是 FP32 的 8 倍，NVIDIA 官方博客 TensorFloat-32 in the A100 GPU Accelerates AI Training, HPC up to 20x 给出的定位就是”用可接受的精度损失换 8-20 倍吞吐”。PyTorch 在 Ampere 上默认把矩阵乘切到 TF32（可手动关闭）；到了 Hopper 时代，BF16 成为张量核心的主打格式，TF32 的使用场景进一步收缩。

FP8：E4M3 与 E5M2 的明确分工#

FP8 是 2022 年由 NVIDIA、ARM、Intel 联合提出的 8 位浮点格式（论文 FP8 Formats for Deep Learning），包含两种编码：E4M3（4 位指数、3 位尾数）和 E5M2（5 位指数、2 位尾数）。它们的设计哲学是把 FP16 的那 11 位”一半给指数、一半给尾数”的两种折中再压缩一次：

• E4M3：范围 $[2^{-9}, 448]$，有效数字 4 位。精度优先，用于前向传播（权重和激活）。为了挤范围，它放弃了 Infinity 表示——指数位全 1 时不再表示无穷，而是让最大正常值扩展到 $1.75 \times 2^8 = 448$（尾数 110），把全 1 尾数 111 留给 NaN。这样动态范围比”按 IEEE 754 规则”多出 1 档。
• E5M2：范围 $[2^{-16}, 57344]$，有效数字 3 位。范围优先，完全遵循 IEEE 754 特殊值约定（有 $\pm\infty$、NaN），用于反向传播（梯度、误差）——因为梯度动态范围大、对精度相对不敏感。

这种分工和 BF16/FP16 的分工逻辑完全一致：计算路径上谁更怕溢出，谁就多拿指数位。

论文的 GPT-3 实验直接验证了这一点：训练损失曲线（perplexity）在不同规模（350M 到 175B 参数）下，FP8 与 FP16 基线几乎完全重合，且超参数一行未改。

有趣的是”哪一个数用 E4M3、指数偏置取多少”本身是个需要实验的数值问题。论文对 1.3B GPT-3 做了系统的指数偏置扫描——把训练好的 BF16 模型权重和激活 cast 成 E4M3，偏置从 0 到 15 逐个试，困惑度差异肉眼可见：

这张图传达的信息是：8 位格式里每一比特都值钱，格式的细微设计（偏置、特殊值）能带来可测的精度差，所以量化工具链都把它们做成可调参数。

三种格式布局放一起看#

Google 的 BF16 论文里有一张经典的格式对比图，把 FP32、BF16、FP16 的位布局和混合精度数据流画在一起：主权重存 FP32，前向/反向计算用 BF16，中间结果在哪一步转换都标得清清楚楚——这正是推理/训练框架里”精度分层”的标准结构（主副本高精度、计算副本低精度、累积再回到高精度）：

softmax 的经典溢出问题#

注意力里的 softmax 是 $\mathrm{softmax}(x)_i = e^{x_i} / \sum_j e^{x_j}$。问题出在指数上：$e^{88} \approx 1.65 \times 10^{38}$ 已经逼近 FP32 最大值 $3.4 \times 10^{38}$，$e^{89}$ 直接溢出成 +Inf。而注意力分数 $QK^\top / \sqrt{d}$ 的取值动辄几十上百（点积方差是 $d$，除以 $\sqrt{d}$ 后方差为 1，但长序列、高维下最大 logit 经常到 30-60，极端情况轻松超过 88）。如果某一行里有 logit $x_i = 100$，那 $e^{100}$ 溢出，整行 softmax 变成 NaN，训练直接崩。

标准解法：减去最大值#

教科书解法是减最大值：令 $m = \max_j x_j$，则

因为分子分母同乘 $e^{-m}$，数学上完全等价，数值上每个指数项 $e^{x_j - m} \in (0, 1]$，绝不会溢出。这是你之前就理解的部分，这里把它作为”数值稳定性重构”的第一个例子：同样的数学对象，不同的计算顺序，数值行为天差地别。

但这个标准做法要求先看到整行数据才能算出 $m$，也就要求先完整算出 $S = QK^\top$ 的一整行、存下来，再扫第二遍。对注意力来说这意味着 $N \times N$ 的中间矩阵要落进 HBM——这正是 FlashAttention 想要消灭的 IO。

分块场景：为什么不能”等最大值”#

FlashAttention 的思路是分块流式计算：一次只把 $Q$ 的一个 block 和 $K$、$V$ 的各个 block 拿进 SRAM，算完 $P \times V$ 的部分和就扔，全程不落盘。问题是：处理第 $t$ 个 block 时，后面 block 里的更大 logit 还没出现——你拿不到全局最大值 $m$，而如果每个 block 各自用”自己的局部最大值”归一化，各部分和的尺度就不同，直接相加是错的。

举个具体例子：$x = [1, 2, 100, 3]$ 分成两块 $[1,2]$ 和 $[100,3]$。第一块归一化后 $[e^{-1}, e^0]/Z_1$，第二块是 $[e^0, e^{-97}]/Z_2$，两者的 $Z$ 一个是 $1.37$、一个是 $1.0000\cdots$——不能相加。必须有一套机制在流式处理过程中持续维护归一化状态，这就是在线 softmax（online softmax）。

在线 softmax：running max + 重缩放#

在线 softmax 算法最早由 NVIDIA 的 Milakov 和 Gimelshein 在 2018 年给出（论文 Online normalizer calculation for softmax，后来成为 FlashAttention 系列的标准组件）。它维护两个累计量：运行最大值 $m_t$ 和运行归一化常数 $l_t = \sum_{j \le t} e^{x_j - m_t}$（注意 $l_t$ 是以”当前最大值 $m_t$“为基准的，这个细节是算法核心）。

处理第 $t$ 个 block $B_t$ 时：

1. 更新最大值：$m_t = \max(m_{t-1},\ m_{B_t})$，其中 $m_{B_t}$ 是本 block 内的最大值；
2. 重缩放旧的归一化常数：因为基准变了，旧项 $e^{x_j - m_{t-1}}$ 要乘上 $e^{m_{t-1} - m_t}$ 才能变成以 $m_t$ 为基准：

3. 重缩放旧输出：输出累加器 $O$（未归一化的加权和）同样整体乘 $e^{m_{t-1} - m_t}$，再加上新 block 的部分和：

全部 block 处理完后，最终答案 $O_T / l_T$。用归纳法可以证明每一步都维持不变式”$O_t / l_t$ = 前 $t$ 个 block 的精确 softmax 加权和”：当 $m$ 不变时只是普通累加，当 $m$ 增大时全体旧项被同一个标量重缩放，比例关系不变。这个算法在数学上是精确的（不引入任何近似），它做的只是把”除法”推迟到最后一步——中间所有项都控制在 $[0,1]$ 量级，这正是它数值稳定的原因。

一个完整的数值例子#

还是 $x = [1, 2, 100, 3]$，$v = [v_1, v_2, v_3, v_4]$，块 1 是 $[1, 2]$，块 2 是 $[100, 3]$。

第 1 块：$m_1 = 2$，$l_1 = e^{-1} + e^{0} \approx 1.3679$，$O_1 = e^{-1}v_1 + e^{0}v_2$。

第 2 块：$m_{B_2} = 100 > m_1 = 2$，于是 $m_2 = 100$。重缩放：

最终结果 $O_2 / l_2 = e^{-99}v_1 + e^{-98}v_2 + v_3 + e^{-97}v_4$（约等于），与一次性算全局最大值的结果逐位一致——$e^{-98}$ 之后那个 $3.7 \times 10^{-43}$ 的旧项在 FP32 里不是零，只是小得看不见，而它在真实 softmax 里本来就只有 $10^{-43}$ 量级，丢弃它对最终结果的影响小于一个 ULP。

注意第 2 步里 $l_1 \cdot e^{2-100}$ 在 FP32 中下溢风险：$e^{-98} \approx 3.7 \times 10^{-43} \ll 1.4 \times 10^{-45}$？不，$3.7 \times 10^{-43}$ 大于 FP32 次正规数下限 $1.4 \times 10^{-45}$，乘上 $1.3679$ 后约 $5 \times 10^{-43}$，仍然可表示（次正规数）。但如果序列再长、旧项再多，这个量级会继续往下掉，最终在 FP32 里下溢为 0——而它在数学上本来就可忽略，所以无害。这里值得强调：在线 softmax 的下溢只会发生在”本来就该被忽略”的项上，不会发生在”正在累积的重要项”上，这正是”把中间值控制在安全区间”策略的威力。

为什么归一化常数 $l$ 必须用 FP32 累积#

一个常被忽视的细节：$l_t = \sum e^{x_j - m_t}$ 的每一项都在 $(0, 1]$，所以 $1 \le l_t \le N$（序列长度）。当 $N > 65504$ 时，$l$ 在 FP16 里必然溢出（65504 是 FP16 最大值）——而 65504 token 的上下文在今天已经是家常便饭。即使没溢出，$l$ 的误差也会随项数线性累积（所有项同号，前面分析过这是最坏情形）。所以 FlashAttention 的 softmax 统计量（$m$、$l$）一律用 FP32 维护，只有 $QK^\top$ 的矩阵乘本身可以用低精度。这是”哪里用低精度、哪里必须保精度”的一个教科书式案例。

与 FlashAttention 的衔接#

你已经在 FlashAttention-4 那篇文章里见过”条件 softmax 重缩放”（conditional rescaling）——那就是这套在线 softmax 机制在 Blackwell 上的工程化：多数 block 处理时 $m$ 根本不更新（重缩放因子 $e^{m_{t-1} - m_t} = 1$），于是 FA4 把”是否真的需要重缩放”做成运行时条件，大多数情况下跳过这步标量乘，把省下的指令周期让给矩阵乘。而 FlashAttention-1（2022）正是靠这套在线机制把 $N \times N$ 的注意力矩阵从 HBM 里解放出来，在 A100 上拿到 2-4 倍加速；FlashAttention-2 进一步优化并行策略后，相对标准 PyTorch 注意力在长序列上的加速比达到接近 10 倍：

数值上”减最大值 + 重缩放”看起来只是一个小技巧，但它直接决定了注意力能不能分块、能不能上 GPU 的 SRAM、能不能用低精度——数值稳定性不是锦上添花，而是整个推理优化栈的地基。

注意力的精度敏感点：指数放大与长尾#

再补一个为什么注意力对数值特别敏感的原因。softmax 的指数函数把输入差异放大：输入相差 $\Delta$，输出概率比相差 $e^{\Delta}$。logit 差 10，概率就差 22026 倍。这意味着注意力分布是高度长尾的：绝大多数概率集中在少数 token 上，其余 token 的概率小到 $10^{-4}$ 甚至更低。低精度格式对”小概率”的相对误差很差：一个 $10^{-4}$ 的概率在 E4M3（$\epsilon = 0.125$）里只能表示成 0 或 $1.25 \times 10^{-4}$ 之类的粗糙值，相对误差几十个百分点。幸好输出 $O = \sum p_i v_i$ 对这些小项不敏感（它们贡献本来就小），这也是为什么注意力可以用低精度矩阵乘、但 softmax 的统计量必须保精度的深层原因——误差必须落在”不影响结果”的地方。

GEMM 的累积器：为什么”算完再存”比”存完再算”准#

矩阵乘 $C = A \times B$ 中，每个输出元素是 $K$ 项乘积的和（$K$ 是收缩维度）。如果每乘一次就把部分和写成 FP16，误差按 $\sqrt{K}$（随机游走）到 $K$（最坏同号）增长；如果全程在 FP32 累积器里累加，最后才写回低精度，误差主要来自最后一次舍入。这就是”累积器精度 > 输入精度”原则：NVIDIA 张量核心的矩阵乘指令（如 HMMA）内置 FP32 累加器，输入可以是 FP16/BF16/FP8/TF32，累积一律 FP32。DeepSeek-V3 更进一步，在 FP8 GEMM 中每累积 128 个元素就把部分和提升到 CUDA Core（FP32 单元）做一次高精度累加，论文 Figure 7(b) 展示的就是这个”每 128 元素提升”的机制：

注意这张图的左半部分（a）画的是另一个问题——离群值。同一张量里如果有少数元素比其他元素大 100 倍（比如注意力投影的某些通道），按全局最大值选缩放因子会让其余元素的有效精度再掉 7 位。DeepSeek-V3 的解法是把缩放因子从 per-tensor 细化到 per-tile/per-block（小到 128 个元素一组），让每一组内幅度相近，这就是”细粒度量化”（fine-grained quantization）。

FP8 为什么必须配缩放因子#

FP8 的 E4M3 范围只有 $[2^{-9} \approx 0.002, 448]$。LLM 的权重通常在 $\pm 0.1$ 量级、激活在 $\pm 10$ 量级，直接 cast 会同时发生灾难性的下溢（权重小值变 0）和上溢（激活大值变 Inf）。所以 FP8 使用前必须做缩放（scaling）：先求缩放因子 $s$，把数据整体放大到格式的满量程，再量化：

$s$ 的取法从粗到细有几种粒度：per-tensor（整个张量一个 $s$，最简单，但被离群值拖累）、per-channel/per-token（权重按输出通道、激活按 token 各给一个 $s$，缓解离群值，代价是反量化多一次乘法）、per-block（把通道进一步切块，每块一个 $s$，精度最好但元数据更多）。缩放因子本身也要存成高精度（FP32），并在计算时融合进 GEMM（$A \cdot B$ 的缩放因子可以合并成一个标量乘）。DeepSeek-V3 的 FP8 框架图把这条链画得很清楚：所有 Linear 算子输入输出都走 FP8，缩放因子 FP32 并行流动，累积在 FP32：

671B 参数用 FP8 训练：损失曲线几乎重合#

DeepSeek-V3（2024 年 12 月，论文）是 FP8 大规模应用的标志性案例：671B 参数的 MoE 模型全程 FP8 训练。它的数值方案值得逐条拆解，因为每一条都是本节讲过的原则的实例化：

• 前向全用 E4M3（不采用”前向 E4M3、反向 E5M2”的常规分工），理由是前向的数值范围可控，E4M3 的精度收益更值——这依赖下面两条兜底；
• 细粒度 tile/block 缩放（Figure 7a）：把缩放粒度做细，把离群值的影响关进笼子；
• 每 128 个元素提升到 CUDA Core 做 FP32 累积（Figure 7b）：累积器保持 FP32 精度；
• 结果：训练损失曲线与 BF16 基线在 EMA 平滑后几乎重合，直到 2.8 万亿 token：

“损失曲线重合”对工程意味着什么？FP8 相比 BF16 把权重/激活内存砍半（对 671B 的模型这是几百 GB 的差别），同时张量核心吞吐翻倍（H100 上 FP8 是 BF16 的 2 倍）。数值格式的每一次降位，都是在”精度 - 内存 - 速度”三角上做一次显式的权衡——这也是为什么你读到的每一篇量化/加速论文，本质上都是一篇数值稳定性论文。

推理侧的其它数值细节#

• RMSNorm 的 eps：x / sqrt(mean(x²) + eps) 里的 eps（通常 1e-5 或 1e-6）不止是防除零，还避免方差极小层的输出被放大到溢出——RMSNorm 的平方和均值用 FP32 累积也是同一套”统计量保精度”逻辑。
• RoPE 的大角度问题：位置 $t$ 很大时旋转角 $t\theta$ 也很大，FP32 里 sin/cos 的精度随角度增大而下降，部分实现用 FP64 计算角度再降回 FP32，或直接查表。这属于”表示精度不足导致结果漂移”的典型场景。
• 整数 vs 浮点量化：INT8 是均匀量化（固定步长），对幅度小的值相对误差极差；浮点格式是相对误差恒定。所以 FP8 天然比 INT8 更适合”数值动态范围大”的 LLM 激活——这也是 2024 年后 FP8 推理（W8A8）逐渐取代 INT8 的数值原因。你之前在 XFP 那篇里看到的码本量化（codebook）走的是另一条路：用多个缩放码本拟合不同幅度的子分布，本质也是”把不均匀的数据用多段均匀/浮点尺度去拟合”。