精度与量化

FP8 是 DeepSeek V3 的三大贡献之一，FP4 权重已经出现在更新的模型上。 所有人都在往下压精度。

但「压精度」不是一个开关，它是一整套决策：

• 哪些张量可以压、哪些必须保住？
• 缩放因子怎么算？per-tensor 还是 per-block？
• 训练时压和推理时压，是完全不同的两件事
• 压过头会怎样？—— 不是「精度略降」，是收敛直接毁掉。

目标：看到一个量化方案，能说出它在哪里省、风险在哪、 以及怎么验证它没把模型搞坏。

• 神经网络对单个数的精度并不敏感 —— 它敏感的是统计上的分布
• 训练本身就带噪声（随机初始化、数据顺序、dropout）， 量化误差在很多地方可以被当成又一种噪声
• 但有些地方不行 —— 那些地方就是这一讲要找出来的

两个收益，性质不同（这一条要讲透，否则后面分不清训练和推理）：

• 算力：低精度矩阵乘吞吐更高 —— 受益的是算力受限的场景（训练、prefill）
• 访存：读写的字节直接减半 —— ⭐ 受益的是带宽受限的场景（decode）

所以量化在推理侧的收益比训练侧直接得多： decode 读一遍权重就是全部开销，字节减半近似就是快一倍。

要立的一条主线：范围（指数位）比精度（尾数位）重要。

• BF16 为什么赢了 FP16：它牺牲尾数保住指数位 ——  动态范围跟 FP32 一样宽，只是精度粗。 溢出是致命的，精度粗一点只是噪声。
• ⭐⭐ 而这条规律会原样再来一次：MXFP4 的 scale 选了纯指数的 E8M0（范围极大、精度极差），NVFP4 选了 E4M3 （有尾数、范围只到 448）—— 同一个取舍，换了个位置又演一遍。 见 §3。
• 还要讲：subnormal 与 flush-to-zero；累加为什么通常仍是 FP32

这一节是整讲的核心。从一个问题往下长：四个 bit 只能表示八个数， 这怎么可能算得准？

先把这一段的问题问出来：E2M1 一共四个 bit， 能表示的绝对值只有八个 —— 这怎么可能算得准？

答案不在那四个 bit 上，在外面套的那层 scale。 四个 bit 只负责「形状」，量级由 scale 给。 所以真正决定它能不能用的，是那个 scale 有多密。

密到什么程度算够？拿一个真尺寸走一遍 ——  设某层权重是 7,168 × 18,432。

• per-tensor：整块一个 scale ——  一亿三千多万个数共用一个数。
• per-channel（按行／列）：一行 7,168 个数共用一个 ——  INT8 时代常见的一档。
• NVFP4：沿收缩那一维每 16 个数一个 ——  那 7,168 会被切成 448 段。

⚠️ 「一个 tensor」是整块矩阵，不是它的某一条边 ——  per-tensor 不等于「按 7,168 量化」，后者是 per-channel。

上一代 FP8 就是最粗的那一档：整张量一个 scale， 而且还得靠软件盯着历史极值去猜它。 一个离群值就把整张量的 scale 拽跑，剩下那些小数全被压进同一个格子。 FP8 有 256 个可表示值，还忍得住；FP4 只有八个，这么干直接崩。

所以 Blackwell 干的事只有一件：把共用 scale 的范围从「整张量」缩到「十六个数」。 这一小撮就叫一个 micro-tensor。

⚠️ 这件事有四个名字，别以为是四件事： NVIDIA 叫 micro-tensor scaling，OCP 标准叫 microscaling（MX） （MXFP4 的 MX 就是它），写 kernel 的人叫 block scaling， 本课叫块量化。

⚠️ 这里有个真的命名碰撞，中文尤其容易踩： 「块」在日常语感里像是二维方块， 于是很自然会以为块量化是「横着几个、竖着几个」拿一小块。

但在这套术语里，「块」不蕴含二维。 NVIDIA 自己给 block quantization 的定义就是沿某一个轴按固定长度分块 —— 文档里连公式带例子：取一个轴当 blocking axis、给一个 block size， 另一个轴不分。所以一维的那 16 个，厂商就管它叫 block。

同一个形状还有别的名字：sub-channel、group-wise —— GPTQ／AWQ 那边的「group size 128」说的就是它； Google 讲 Ironwood 调优的公开文章里也写作 subchannel activation quantization。 它们和 block scaling 指的是同一件事，只是出身不同的社区。

真正的二维块量化是另一回事，而且一般会把维度写出来 ——  比如 DeepSeek V3 的 FP8 权重是 128×128 的二维块。 看到「块量化」而没写维度，默认它是一维的。

两个格式的数值部分完全一样，差别全在 scale 上：

• MXFP4：每 32 个数一个 E8M0 的 scale ——  纯指数没有尾数，所以倍数只能是 2 的整数次幂。
• NVFP4：每 16 个数一个 E4M3（就是 FP8）——  有尾数，不再卡在 2 的幂上；外面再套一个整张量共用的 FP32，两级。

代价是平均位宽：NVFP4 是 4 ＋ 8/16 ＝ 4.5 bit， MXFP4 是 4 ＋ 8/32 ＝ 4.25 bit。

⚠️ 那两级 scale 各是什么类型？—— 差别很大，而且是算出来的，不是随便定的。

• 块内那一级：8 bit。 NVFP4 用 UE4M3（不带符号位，4 指数 3 尾数）； MXFP4 用 E8M0（纯指数，连尾数都没有）。 都不是 FP32，也不是 BF16。
• 整张量那一级：真的是 FP32，而且真的是 per-tensor —— 不是 per-channel。 官方原话就是「一个额外的、单一的 per-tensor FP32 缩放因子」。 一整张量只有这一个数。

为什么块内那级只能给 8 bit —— 因为它要摊到每个值头上。 NVFP4 每 16 个值配一个 scale，8 ÷ 16 ＝ 每个值多摊 0.5 bit。
换成 BF16 就是 16 ÷ 16 ＝ 每值多摊 1 bit，4 bit 的格式变成 5 bit； 换成 FP32 就是多摊 2 bit，直接变 6 bit ——  省下来的那一半全吐回去了，还不如直接用 FP8。

而 per-tensor 那一级不用摊，一整张量就一个数， 占多少位都约等于零 —— 所以它可以大方地用 FP32。

⭐ 这就是这套格式的设计逻辑：精度砸在不用摊的那一级上。

⚠️ 那内层外层到底怎么分工？——  外层管范围，内层管局部量级。 两级不是「更高级」，是为了补一个坑。

坑在这里：NVIDIA 想让块内 scale 带尾数（比 2 的幂更准）， 于是挑了 E4M3。 但 E4M3 自己能表示的最大值只到 448 ——  而一个真实张量里，不同块之间的量级可以差出好几个数量级， 448 那点范围根本盖不住。

所以才有外层那一级：先用一个 FP32 的全局 scale 把整张量搬进内层能表示的区间，剩下的局部差异再交给每 16 个数的 E4M3。 NVIDIA 自己对这一级的定性是「防溢出」，不是「提精度」。

范围能算出来：一个值最终 ＝ 全局 FP32 × 块内 E4M3 × 那个 E2M1。 后两级叠满就是 448 × 6 ＝ 2,688 ——  所以全局那一级大致取「整张量里最大那个块的 amax ÷ 2,688」， 正好把最大的块顶到不溢出的位置。

⭐ 反过来就能看懂 MXFP4 为什么只有一级： E8M0 是纯指数、8 位全给幂次，范围大得离谱 —— 根本不存在「盖不住」的问题。 它不需要外层，代价是那个 scale 只能是 2 的整数次幂，粗。
一句话：NVFP4 拿「多一级」换「scale 更准」，MXFP4 拿「scale 更粗」换「结构更简单」。

⚠️ 那机器怎么知道自己拿到的是哪一种？—— 写在指令里，三处限定符。

• kind：NVFP4 是 kind::mxf4nvf4， MXFP4 是 kind::mxf4。
• scale_vec：NVFP4 是 4X，MXFP4 是 2X —— 这个倍数不用查也能对上：一条指令 K＝64， 64 ÷ 16 ＝ 4 个 scale，64 ÷ 32 ＝ 2 个。
• stype（scale 的类型）：NVFP4 是 ue4m3， MXFP4 是 ue8m0。

⚠️ 外层那个 per-tensor FP32 不在指令里 ——  它是软件在量化和收尾时自己乘进去的，硬件只认块内那一级。

⚠️ 那 16 个数是怎么排的？—— 是一条线，不是一个方块。

很自然会猜它是「横 4 竖 4」的小方块。不是。 它是一行里沿收缩维连续的 16 个数，完全不跨行。

不用信转述，直接读那条指令的分片形状： NVFP4 那条 mma.sync 吃进的 A 片是 16×64（M×K）， 而它的 scale 片 SFA 是 16×4。 64 ÷ 16 ＝ 4 —— 一行 64 个 K 切成 4 段，每段一个 scale。 B 那边同理：B 片 64×8，SFB 是 4×8。

所以落到激活上，一个 block ＝ 同一个 token 的连续 16 个特征。 它一个 token 都不跨。

那为什么非得沿收缩维？ 因为 scale 必须能从求和号里提出来。 沿 K 切成一段一段，每段内 sA、sB 都是常数，那一段的点积就是 sA·sB·Σ(a·b) ——  一段只要一次额外乘法。
换成「每个通道一个 scale」就提不出来了 —— 每一项的系数都不同，求和号里再也拆不开。 所以「沿 K 分段」是「仍然能提出来」的前提下能做到的最细粒度。

⭐ 最要紧的一点：它不需要「相邻的数值相近」这个假设。 NVIDIA 给的理由从头到尾只有一条 ——  更多次机会去贴合局部的动态范围。
要的只是「这 16 个数的量级差得别太远」，不是「它们彼此相似」。 而这是个统计性质，跟相不相关无关： 十六个数里出一个离群值的概率，远低于一亿三千万个数里出一个； 真出了，被拖累的也只有那十五个邻居。

⚠️ 但要小心，别把「硬件是一维」读成「二维块缩放没人用」。

NVIDIA 自己的 NVFP4 训练 recipe，权重那一侧就是 16×16 的二维块 —— 官方文档原话是「一个 scaling factor 由一个 16×16 的二维块共享」， 并说这跟 DeepSeek V3 训练里的权重量化是同一路子，只是粒度细得多。

这两件事不矛盾，它们在不同层： 硬件那一层永远是一维 16；二维是叠在它上面的一层约束 —— 算出一个 16×16 块的 scale 之后，让那 16 行各自的 1×16 共用这一个值。 （这一句不是推测：论文原文就是「2D block scales are replicated for each of the 1×16 blocks when being passed into Tensor Cores」。）

先回到专题四那张 16 字节的表，从精度角度重看一遍：

• 前向／反向用 bf16 —— 省算力省带宽
• 主权重必须用 fp32 —— ⚠️ bf16 只有 7 位尾数， 训练后期的小更新量加上去会被直接舍掉，等于没更新
• 优化器状态、累加 —— 同理，fp32

⭐ 那量化那一步，具体是怎么算的？

拿一个 16 个数的块走一遍（推理侧、NVFP4）：

• ① 先在整张量上取一个全局 scale α（FP32）。 取法大致是「最大那个块的绝对值最大值 ÷ 2,688」 —— 2,688 ＝ 448 × 6，就是后两级叠满能表示的上限。
• ② 在这 16 个数里取绝对值最大值 amax。
• ③ 块内 scale Δ ＝ amax ÷ 6 ÷ α， 再把它round 成 UE4M3 存起来。 （除以 6 是因为 E2M1 最大只能表示 6；除以 α 是把它挪进 E4M3 的量程。）
• ④ 每个元素 q ＝ round( x ÷ (α · Δ) )， 结果落进那八个可表示值里，存成 4 bit。 权重和激活用就近偶数舍入，梯度用随机舍入。

反量化就是把这条路倒着走：x ≈ α · Δ · q。 其中 Δ · q 那一步是 Tensor Core 在 MMA 里顺手做的， α 那一步是软件在 FP32 输出上乘回去的。 （所以硬件只认两级里的一级 —— 这一条前面说过，这里对上了。）

⭐ 再补一个图景，不然下一段听不懂：一个线性层其实是三个矩阵乘。

• 前向（Fprop）：y ＝ w · x
• 反向之一（Dgrad）：∂x ＝ wᵀ · ∂y —— 算给上一层的梯度。
• 反向之二（Wgrad）：∂w ＝ ∂y · xᵀ —— 算这一层权重自己的梯度。

三个都在 FP4 上算，输出是 BF16 或 FP32。 注意第二个：它用的是 wᵀ —— 同一个权重，被转置着又用了一次。 下一段整段就是在处理这一件事。

⭐ 那 16×16 到底怎么回事？—— 这是训练侧的事，而且理由不是「更准」。

先把配方摆出来（出自 NVIDIA 那篇 NVFP4 预训练论文）：
权重用 16×16 的二维块， 激活和梯度仍然用 1×16 的一维块。 （那 16×16 是「16 个输入通道 × 16 个输出通道」。）

为什么权重要特殊照顾？—— 因为反向传播会把它转置。

缩放必须沿点积那一维，这条前面说过。 可前向算的是 y ＝ w·x，反向算的是 ∂x ＝ wᵀ·∂y —— 点积那一维换了。
于是同一个权重，在前向被沿行切块量化，在反向被沿列切块量化， 得到两份不一样的量化结果。

⛔ 这不是「精度差一点」，这是链式法则被打破。 论文说得很直白：反向传播算的梯度，对应的已经不是前向那个函数了 —— wfprop ≠ wbprop， 你在给另一个函数求导。

16×16 就是来堵这个的：一个正方形的块，转置之后还是同一个块。 前向反向拿到的是同一份量化权重，链式法则重新成立。

⭐ 最反直觉的一点：16×16 是 256 个数共用一个 scale， 比 1×16 粗十六倍。 它更粗，却更好 —— 论文的原话就是「尽管块粒度更大， 2D 仍然改善了 loss」。说明这里赢的不是精度，是一致性。

那硬件在这里干了什么？—— 什么都没多干。

Tensor Core 只认沿点积维的一维 scale，这是硬性约束 —— 论文自己也这么写。
所以 16×16 是这么落地的：软件按 16×16 算出一个 scale， 然后把它复制十六份，当作那 16 行各自的 1×16 scale 喂进 Tensor Core。 硬件从头到尾不知道有「二维」这回事，它看到的还是十六个一维块 ——  只不过那十六个 scale 的值恰好一样。

⚠️ 所以二维不省存储：喂给硬件的 scale 张量大小跟一维时一模一样， 省的不是空间，买的是「转置前后是同一份」。外层那个 per-tensor FP32 照旧。

⚠️ 论文里确实出现过 4，但那是另一个东西。 它比较的 d ＝ 4 / 16 / 128 是随机 Hadamard 变换的矩阵尺寸 —— 用来把离群值打散的那一步，跟块大小不是一回事。 结论是 d ＝ 16 比 4 收敛好、跟 128 差不多，所以选了 16。 两个 16 同值不同物，最容易混。

配方里还有两条跟量化方式直接相关的： 梯度用随机舍入（去掉量化偏置）， 权重和激活用就近偶数舍入； 数值敏感的层整层留在高精度。 （另外论文发现：训练后段把前向切回 BF16 能把 12B 模型的相对误差从 1.5% 拉回 0.5%，而切反向没有同样的收益。）

同一路子的还有 DeepSeek V3 的 FP8 权重 —— 128×128 的二维块。 所以判据要改一句：看到「块量化」而没写维度，说的多半是硬件那层的一维； 但训练侧的 recipe 很可能在上面又叠了二维，得看具体 recipe。

这一节接专题二。 「硬件支持 FP4」这句话，落到硅上到底是什么？

所以 HBM 里放的是 FP4，不是 FP8 ——  打包好的 E2M1 加那一串 scale，省下的一半就是这么省的。

那「硬件支持 FP4」到底加了什么？—— 没有协处理器。 加的是三样，分布在三个不同的地方：

• ① 乘法阵列本身认 4 bit。第五代 Tensor Core 原生支持 FP4 与 FP6 —— 这是「能算」。
• ② MMA 数据通路里内嵌一级「乘 scale」。 这就是反量化 —— 不是独立部件，是流水线里的一步。
• ③ 通用核那边新增转换指令（把 FP32 打包成 E2M1 那类）。 这是「量化」—— 跑在普通 CUDA Core 上，不是专用引擎。

有一个反例能把 ① 和 ③ 是两块独立硬件钉死： 桌面那颗 GB10 有 FP4 的 Tensor Core 指令，却缺那条转换指令 —— 能拿 FP4 算矩阵乘，却没法把数转成 FP4。 同一颗芯片上，一个有一个没有。 （出自 NVIDIA 开发者论坛的实测报告，不是官方规格书。）

于是分工很清楚：硬件负责「用」scale，软件负责「算」scale。 权重是离线量化好、直接以 FP4 存进去的；激活是运行时由 kernel 在上一步收尾里算的。

硬件那条乘法长这样：指令是 tcgen05.mma 的 block-scale 变体 （B200 走这条）和 warp 级的 mma.sync…block_scale （消费级那颗 die 走这条）—— 这正是上面那道门槛的来处。
A、B 是 FP4，两串 scale 叫 SFA／SFB， Tensor Core 一边读 E2M1 一边把 scale 乘进去，累加仍是 FP32。

⚠️ scale 不是「放进 Tensor Core 里」，是放进 TMEM。 TMEM 是 SM 上的一块内存 —— 每 SM 256 KiB， 128 lane × 512 column、每格 32 bit，路径是 显存 → 共享内存 → TMEM。 （⭐ lane ＝ TMEM 的「行」，一条 lane 对应一个线程的落点 —— 所以下面那句「一个 warp 只够得着 32 条 lane」说的就是 warp 的 32 个线程。） 它装的是累加器和 scale，不是 FP4 操作数本身（那些从共享内存流进来）。

⭐ 「乘 scale 不额外消耗算力」这句对 ——  那一乘是流水线里的一级，不占发射槽，也不吃 FP4 的峰值。 但「不花钱」是另一回事，代价有三笔，只是都不算在 FLOPs 头上：

• ① 占 TMEM。一个 128×256 的 FP32 累加器就是 128 × 256 × 4 B ＝ 128 KiB，那 256 KiB 一下去掉一半， scale 得挤在剩下那半里还不能撞 bank。 真正卡住你能开多大 tile 的，往往是这一笔。
• ② 多一路搬运。那串 scale 自己占 TMA 管线和共享内存， 而且必须摆成一种交错布局 —— 天真按 K 排好的那种得先重排。
• ③ 量化那一头是真花时间的。算每块的最大值、压成 E2M1， 跑在通用核上 —— 不进 MMA 的账，但进墙钟的账。 只比 MMA 就会把这一段漏掉。

连起来，一层的来回是这样一圈：
① 上一层的结果在 TMEM 里是 FP32 累加值 →  ② 收尾时就地量化成 FP4 ＋ scale，写出去 →  ③ 下一层把 FP4 搬进共享内存、scale 搬进 TMEM →  ④ Tensor Core 边读边乘 scale，累加又回到 FP32 → 回到 ①。

⭐ 一句话：存的是 4 bit，算的是 32 bit。 进来那一下的换算焊在 MMA 流水线里，出去那一下 是普通核跑一条新指令。

课后延伸　四块「问到硅上」的细节

下面四块都是把上面那三样往下再追一层， 主线上跳过不影响听懂。第 ② 块最实用 ——  它能解释「为什么我换了 FP4 却没快一倍」。

延伸一：为什么一套硬件能同时吃 NVFP4 和 MXFP4

先回答那个问题本身：两个都支持，而且是同一套单元。

关键在于：scale 不管哪种格式，永远是一个八位的无符号浮点数。 区别只在这八位怎么解释 —— 位宽一样，数据通路就能共用。

操作数格式	一个 scale 盖多少个元素	scale 的类型
mxf8（E5M2 / E4M3）	32	UE8M0
mxf6（E3M2 / E2M3）	32	UE8M0
mxf4（E2M1）＝ MXFP4	32	UE8M0
nvf4（E2M1）＝ NVFP4	16	UE4M3

两种 scale 类型都是无符号八位、都支持 NaN 但不支持无穷。 UE8M0 把八位全部给指数（取值只能是 2^x， −127 ≤ x ≤ 127）；UE4M3 是四位指数三位尾数。

⭐ 这个差别在硬件上是实打实的成本差： 乘一个 2 的整数次幂，等于把浮点数的指数字段加一下 ——  几乎白送。而 UE4M3 带尾数，那是一次真的乘法。 所以 NVFP4 多出来的那点精度，是拿电路换的 —— 这也是为什么 MX 是行业标准、NVFP4 是自家格式。

⭐ 还有一件事挺有意思：硬件比这两个格式加起来更宽。 PTX 里 mxf4nvf4 这一档，块长 16 时 E8M0 和 E4M3 都收， 块长 32 则必须配 E8M0。 也就是说「块 16 ＋ 纯指数 scale」这个组合硬件是支持的， 只是没有哪个标准格式叫这个名字。

⚠️ 延伸二：同样是 FP4，走错指令只有一半速度

这一条是本讲主线最锋利的一个实例。 NVIDIA 的 CUTLASS 文档里，Blackwell 的 MMA 指令是按 kind 分档标速度的（倍率相对 Hopper 的 FP8 Tensor Core）：

指令 kind	能吃什么	速度档	矩阵布局
kind::f8f6f4	f4 / f6 / f8 混着来	2×	四种都行
kind::mxf8f6f4.block_scale	mxf4 / mxf6 / mxf8	2×	四种都行
kind::mxf4.block_scale	只吃 mxf4	4×	只有 TN
kind::mxf4nvf4.block_scale	mxf4 或 nvf4	4×	只有 TN

TN ＝ A 行优先、B 列优先。4× 那两条只认这一种排法。

⭐ 所以「B200 的 FP4 是 FP8 的两倍」这句话有两个隐含前提： 你走的是专用那条 kind，而且矩阵摆的是 TN。

走那条什么都能吃的混合指令会怎么样？ 数据还是四位的、显存还是省一半、精度行为也一样 —— 只有矩阵乘的速度掉回 FP8 那一档。 而这件事不会报错，也不会在任何日志里说一句。

⛔ 这就是这一讲第 6 节说的「沉默的失败」在性能上的版本： 规格表给的是上限，门槛写在指令表里 —— 而两者之间隔着一整份你没读过的文档。

延伸三：标称的 FP4 算力里，把 scale 和反量化算进去了吗

没有。标称只数乘加，一次乘加记两个 FLOP。

有一条很干净的旁证能把这件事钉死 —— 不用查任何白皮书，用两个格式互相对照就行：

• NVFP4 是每 16 个数一个 scale，MXFP4 是每 32 个数一个 scale —— scale 的数量差整整一倍。
• 而 B200 的 FP4 峰值对外只报一个数，没有按格式分开。
• ⭐ 如果 scale 进账，这两个格式的峰值必然不同。 它们相同，所以 scale 不在账里。

那「反量化算不算一次 FLOP」？—— 这个问题问不出来，因为 没有一个单独的反量化步骤可数。 它就是 MMA 流水线里的一级。你要么用 block-scale 的 MMA （反量化长在里面，拆不出来），要么自己先把 FP4 升成 BF16 再算 —— 那走的是 BF16 的峰值，压根不是 FP4 那个数字。

⚠️ 反过来有一笔真该记住的账：量化那一侧是实打实花时间的， 而且一个 FLOP 都不在峰值里。 算每块的最大值、压成 E2M1、生成 scale ——  这些是真实指令，跑在通用核上。 只比 MMA 的峰值，就会把这一整段漏掉。

延伸四：TMEM 为什么会存在 —— 它不只是「一块新内存」

先看它取代了什么。 Hopper 及以前，MMA 的累加器住在寄存器堆里，每个线程拿着结果的一小片。 Blackwell 的 MMA 块子大到累加器会把寄存器吃光，占用率直接被压死。

于是给累加器单开了一块地。CUTLASS 文档的说法是： 这样一来，寄存器就能主要用于调度和 epilogue 那些事。

换来的三样东西，一样比一样重要：

• ① 寄存器压力解除 —— 能开更大的 tile。
• ⭐ ② MMA 可以（而且必须）由单个线程发起。 因为这条指令要用的数据一个字节都不在私有寄存器里， 全在 CTA 共享的内存空间中。 矩阵单元跟 warp 调度器解耦了 —— 这是三样里最深的一样。
• ③ 两个 SM 能合做一次 MMA（cta_group::2） —— 只有累加器住在公共内存里才做得到。

代价也很具体：这块地要你显式申请、显式释放； 分配的单位是列（必须是 2 的幂、至少 32 列）； 算完还得用专门的指令把结果搬回寄存器才能做后处理； 而且一个 warp 只够得着 32 条 lane（四分之一） —— 所以光是搬结果就得动用一整个 warpgroup。

⭐ 把「分配单位是列」这一条记住，上面第 ① 笔代价就通了： 存 scale 只用到 32 条 lane，但另外 96 条也跟着被占住、别的东西不能用 —— 因为一列的 128 条 lane 是一起给出去的，没法只要四分之一。

量化出问题往往不报错。已经攒下三类，一类比一类隐蔽：

• ① 编译过了、跑起来了、loss 也在降 ——  但降得比该有的慢，或者后期崩
• ② 开关没生效也不报错 —— QAG 那五道锁只开一把就是这样
• ⭐ ③ 链式法则被打破（wfprop ≠ wbprop） —— 它连「开关没生效」都算不上：一切正常，只是你在给另一个函数求导。 见 §4。

所以：每一次精度改动都必须有对照实验，而且要跑足够长才看得出差别。

按「结论 ← 材料」排，不按书目排 ——  要回查的人手里拿的是结论，不是书单。 凡是推的都单独标了。

⭐ 出处台账 —— 上面每一条结论是从哪儿来的

• E2M1 那八个可表示值、两级缩放、16 对 32、E4M3 对 E8M0、 「更多次机会贴合局部动态范围」
  ← NVIDIA 开发者博客 Introducing NVFP4 for Efficient and Accurate Low-Precision Inference
• MX 标准的 k＝32、E8M0
  ← OCP Microscaling 规范及其论文 Microscaling Data Formats for Deep Learning
• 那 16 个是一维沿 K（A 片 16×64、SFA 片 16×4）、UE4M3 这个精确叫法、 NVFP4 那条 PTX 指令串
  ← Colfax Research NVFP4 Blockscaled GEMM on NVIDIA RTX Pro Blackwell GPUs (SM12x) ·「每个 scale 盖 K 方向连续 16 个元素」另见 NVIDIA cuDNN 前端关于 scale factor 布局的文档
• UE8M0／UE4M3 的位分配与取值范围、四种操作数格式对应的块长与 scale 类型、 mxf4nvf4 块 16 可配两种 scale、TMEM 按列分配（2 的幂、≥32 列）、 一个 warp 只够得着 32 条 lane
  ← Colfax Research Hardware-supported Block-scaling with NVIDIA Blackwell GPUs 与 Writing GEMM Kernels Using Tensor Memory 后者也是「累加器搬出寄存器堆是为了把寄存器让给调度和 epilogue」「所有数据都在 CTA 共享空间里，所以这条指令可以且必须由 单线程发起」两句的来处。
• ⭐ 四条 MMA kind 的速度档与布局限制（2× 对 4×、 4× 只支持 TN）
  ← NVIDIA CUTLASS · Blackwell SM100 GEMMs 的 tcgen05.mma 指令表
  ⚠️ 表里标的是相对 Hopper FP8 Tensor Core 的倍率， 不是绝对 TFLOPS。「走混合 kind 就掉回一半」是照这张表读出来的， 不是实测 —— 想确认的话得自己跑两条 kind 对比。
• 「标称 FP4 峰值里不含 scale」
  ← ⚠️ 这一条没有官方声明，是推的，推导链如下：
  NVFP4 每 16 个元素一个 scale、MXFP4 每 32 个一个 （见上面两条的出处）→ 两者 scale 数量差一倍 → 而 B200 的 FP4 峰值 对外只公布一个数、不按格式区分 → 若 scale 计入 FLOP，两个格式的 峰值不可能相同。反证成立，但不是官方口径。
• scale 必须进 TMEM、SFA／SFB 的交错布局、显存→共享内存→TMEM 这条路
  ← Colfax Research Hardware-supported Block-scaling with NVIDIA Blackwell GPUs，以及 NVIDIA 的 tcgen05 MMA 文档
• TMEM 每 SM 256 KiB、128 lane × 512 column × 32 bit
  ← Colfax 的 TMEM 教程；另有公开微基准论文 Microbenchmarking NVIDIA's Blackwell Architecture 独立复核
• 「block quantization ＝ 沿一个轴分块」这个定义（不蕴含二维）
  ←  NVIDIA TensorRT · Quantization Workflows （文档里连 blocking axis、block size 和 scale 形状的公式都给了）
• per-tensor 那一级是「防溢出」不是「提精度」；训练 recipe 的权重用 16×16
  ←  NVIDIA Transformer Engine · Using FP8 and FP4 （原话：an additional single per-tensor FP32 scaling factor to avoid overflows；a single scaling factor is shared by a 2D block of 16x16 elements）
• 16×16 的真正理由（链式法则）、2D 比 1D 粗却更好、 scale 复制成 16 份喂 Tensor Core、Hadamard 的 d＝16、 随机舍入、后段切回 BF16 的收益
  ←  Pretraining Large Language Models with NVFP4（NVIDIA） ·「Tensor Cores require scaling factors along the dot-product dimension」这句硬约束也出自该文
• GB10 有 FP4 的 MMA 却缺转换指令
  ← NVIDIA 开发者论坛上的实测报告 ⚠️ 这是论坛帖，不是官方规格书 ——  本课里它只用来支撑「能算」和「能转」是两块独立硬件这个定性判断， 不用它支撑任何数字。

⚠️ 本段里属于「推的」只有一处，已在原处标明： 「集合通信为什么也适合交给 SparseCore」那一句 —— 那是 3.5 另一半的事，跟 FP4 无关，写在这里只是为了台账完整。 FP4 这一段的每一条都能对上上面某一份材料。

⚠️ 最后澄清一个名字：「Transformer Engine」不是物理部件 ——  它是「支持块缩放的 Tensor Core ＋ 决定怎么选 scale 的软件库」的合称， 芯片版图上找不到它。所谓第二代，差别就是上面那条： 第一代整张量、软件猜；第二代每十六个一个、硬件乘。