日期：2026-08-28
涉及硬件：NVIDIA Blackwell（B200/GB200，数据中心 SM100；RTX 50 系列，消费级 SM120）、Hopper（H100/H200）、Ampere（A100）
涉及格式：FP4（E2M1）、NVFP4、MXFP4、FP8、FP16/BF16
阅读前提：只需要知道「大模型是个很大的神经网络」。二进制、浮点数、量化、显存带宽、tensor core，全部从零讲起，每一步只依赖上一步。
来源说明：英伟达官方博客、PTX 指令集文档、OCP MX 规范里的内容是公开事实，文末统一列参考文献；tensor core 内部电路英伟达从未公开，涉及内部实现的地方文中标「推断」。
1. 开篇：为什么是 NVFP4
这一章只有两件事：三笔真账，和一个疑问。
先看一笔账：DeepSeek-R1 有 6710 亿参数。每个参数用 FP16（16 位浮点数）存，光权重就是 671e9 × 2 字节 ≈ 1.34 TB。一张 B200 显存 192 GB，8 张 DGX B200 一共 1.5 TB——装是装下了，但没剩多少地方放中间结果。
二笔账：模型装下只是开始。推理时每生成一个字，显卡要把参与计算的权重从显存里搬一遍、乘一遍。参数越多，搬得越多；搬运有多快，直接决定出字有多快——这句话第 4 章会用 H200 的真数据验证。
三笔账：2025 年 3 月，英伟达官宣 DGX B200 用 4-bit 精度（FP4）跑 DeepSeek-R1 671B，整机吞吐超过 30000 tokens/s，精度还贴着 16-bit 基线。同年 8 月，他们又发了一组更狠的数据：120 亿参数模型直接用 NVFP4 从头预训练 1 万亿 token，loss 曲线和 16-bit 基线几乎重叠。
「又小又快还不掉精度」，这三样凑齐，行业只剩一个反应：跟进。2025 年下半年起，vLLM、SGLang、TensorRT-LLM、PyTorch 相继把 NVFP4 写进量化支持列表。AMD 和 Intel 则押注同一套 OCP 开放标准里的另一个 4-bit 格式 MXFP4（第 5 章对比）。
但这里有个拧巴的地方：4 个比特一共只有 2⁴ = 16 种组合，一个 4-bit 数最多表示 16 个不同的值。用 16 个值去装 6710 亿个参数——这不是压缩，这是魔法。
这篇文章把魔法拆开讲：16 个数长什么样、裸用为什么报废、NVFP4 的「缩放」怎么盘活它们、显卡硬件到底在哪一步帮了忙、以及最核心的一步——缩放可以被挪到求和符号外面，这是 4 倍吞吐成立的全部前提。
先给答案，下面各章展开。
裸的 4-bit 浮点（E2M1）只有 16 个值，量程 ±6，直接拿来用不行。NVFP4 的命根子是两级缩放：每 16 个元素共享一个 FP8（E4M3）块 scale，整个张量再共享一个 FP32 全局 scale，每个元素实际只花 4.5 bit。
4 倍吞吐不是因为「4-bit 数学更便宜」，而是同一块硅把乘法 lane 切细了 4 倍（推断）。最关键的技巧：块 scale 是块内常数，可以提出点积的求和号，乘法 lane 里流转的永远是原始 4-bit 值。
软件能复刻全部数学（延迟缩放，精度等价），复刻不了电路——自由 CUDA 代码摸不到 4-bit 乘法 lane，吞吐封顶原生的 1/4。老卡跑 4-bit 同理：显存和带宽的收益照拿，算力收益封顶本代 tensor core。
一句话：门是开着的（PTX 公开、CUTLASS 开源），门后只有一条固定传送带。
2. 补课一：从比特到 E2M1 的 16 个值
这一章从零打地基：权重是什么、数在内存里怎么存、E2M1 的 16 个值怎么一个一个推出来。会浮点数的读者可以直接跳到 2.4 看值表。
2.1 权重是什么
神经网络前向传播干的事情，剥掉所有修饰就是一连串的「乘 + 加」：输入的每个数乘上一堆预先存好的数，乘完加起来，再过一道非线性函数，喂给下一层。
那堆「预先存好的数」叫权重。大模型所谓的「大」，就是权重的数量大：70 亿参数的模型，就是 70 亿个这样的小数。后文反复说的「砍位数」，砍的就是每个权重占的字节数。
2.2 比特和二进制
计算机里最小的存储单位是比特（bit），只有 0 和 1 两个状态。8 个比特组成一个字节（byte）。任何数存进内存，最终都是一串 0 和 1。
十进制的 5，二进制写成 101：从右往左分别是 1 个 4、0 个 2、1 个 1。小数也一样，只是底数换成 2 的负指数：0.75 = 0.5 + 0.25 = 2⁻¹ + 2⁻²，写成二进制就是 0.11。
所以一个 0.75 用 FP32 存进内存，长这样：
图 1：0.75 在内存里的样子：32 个比特分三段——1 位符号、8 位指数、23 位尾数。
先不解释后两段，记住一个感觉：一个数在内存里的样子，是「符号 + 指数 + 尾数」三段 0/1 拼出来的。格式不同，就是这三段各分几个比特不同。
2.3 浮点数 = 二进制世界的科学计数法
科学计数法大家见过：光速 3.0×10⁸。任何非零数都能拆成「1 位有效数字 × 10 的某次方」。浮点数（floating point）就是把这套搬到二进制：
数值 = ± 尾数 × 2^指数「浮点」的意思是：指数一变，小数点的位置跟着滑动，像整把尺子平移。一个浮点格式要规定的东西就两件事：总共几个比特，怎么在符号/指数/尾数之间分。
| 格式 | 位分配（符号+指数+尾数） | 最大正数 | 谁在用 |
|---|---|---|---|
| FP32 | 1+8+23 | 约 3.4×10³⁸ | 老一代训练 |
| FP16 | 1+5+10 | 65504 | 训练主流 |
| BF16 | 1+8+7 | 约 3.4×10³⁸ | 大模型训练（范围同 FP32，精度降档） |
| FP8 E4M3 | 1+4+3 | 448 | Hopper/Blackwell 推理、低精度训练 |
| FP8 E5M2 | 1+5+2 | 57344 | 梯度（范围优先） |
| FP4 E2M1 | 1+2+1 | 6 | Blackwell 低精度计算 |
表里有个容易忽略的行：BF16。它是 FP32 直接砍出来的——保留 FP32 的 8 位指数（范围完全一致），尾数砍到 7 位。
大模型训练为什么偏爱它而不是 FP16？因为训练里梯度会漂，FP16 范围只有 ±65504，一漂就溢出；BF16 范围和 FP32 一样宽，丢的只是尾数精度，而神经网络扛得住（第 3 章讲的稀释机制）。
名字里的 E 和 M 是指数（Exponent）和尾数（Mantissa）的位数。E2M1 = 2 位指数 + 1 位尾数，加上 1 位符号，正好 4 个比特——这就是「FP4」的全部。
图 2：四种浮点格式的位布局：E2M1 把符号/指数/尾数砍到 1/2/1，4 位只有 16 个值，量程 ±6。
两段比特各管一件事：指数位管范围（这个数能大/小到什么量级），尾数位管精度（同一个量级内部能切多细）。指数每多 1 位，上限翻一倍量级；尾数每多 1 位，每个量级内部的刻度密度翻倍。砍位数，砍的就是这两样。
2.4 E2M1 的 16 个值
现在把 E2M1 能表示的数全部推出来。指数 2 位，取值 00、01、10、11；尾数 1 位，取值 0 或 1。约定（OCP 规范）：E2M1 的指数偏置是 1，即实际指数 = 指数位 − 1；指数位全 0 时按「次正规数」处理，尾数不再是「1.几」而是「0.几」。
逐组合算一遍，再乘上符号位（0 正 1 负）：
图 3：上表是 8 种无符号组合的逐行推导（指数位 × 尾数位，次正规行高亮），下面一排色块是乘完符号位后的全部 16 个编码——这就是 FP4 能表达的全部。
一共 16 个（0 的两种符号编码都算 0）。这 16 个数就是「FP4 能表达什么」的全部——它们扛得住多大的活，下一章说。
2.5 既然都是 4 个比特，为什么不直接用 INT4
顺理成章的疑问：4 个比特存 16 个整数（-8 到 7），不行吗？
不行，差在刻度分布。INT4 的 16 个刻度是均匀排开的，每格 1。但神经网络的权重分布是「中间密、两头疏」：绝大多数权重挤在 0 附近（±0.1 量级），少数离群值（outlier）冲到 ±几十。
用均匀刻度去量这种分布，0 附近最该密的地方反而最疏——大量刻度浪费在没人待的大数值区间。
浮点的刻度天生「小数值处密、大数值处疏」：指数一动，整把尺子平移一个量级。这和权重分布的脾气正好相配。这就是低精度竞赛里人人选浮点、不选整数的根本原因。
（顺带一句：INT4 并没有死，它走另一条路——均匀刻度配更激进的缩放策略，GPTQ/AWQ 那一批量化方案就是 INT4 阵营。本文主角是浮点这条。）
3. FP4 的 16 个值，为什么不能直接用
这一章回答两个问题：16 个值裸用为什么报废，以及「砍位数」到底损多少、损怎么被兜住。
3.1 裸 E2M1 的三个致命伤
回看 2.4 那 16 个值，三个致命伤一眼就能数出来：上限是 6，超过的数直接溢出，没有「略大于 6」这一档；最小非零值是 0.5，0.1、0.02 存不了，只能舍入到 0 或 0.5；刻度不均匀还稀，1 和 1.5 之间什么都没有，2 和 3 之间也什么都没有。
而真实权重的分布是：绝大多数挤在 ±0.1 的窄带里，少数 outlier 冲到 ±几十。裸 4-bit 怼上去，两头都崩——大的溢出，小的塌成 0。
3.2 28.6% 的账：量化误差到底损在哪
把 FP16 权重「存成」FP4，动作只有两个：找最近的档位、四舍五入。这个动作叫量化（quantization）；反着来——从档位还原回高精度数——叫反量化（dequantization）。
拿一个具体数算损失。设某个权重是 0.7（注意：裸 E2M1 的刻度是 0、0.5、1、1.5、2、3、4、6）。0.7 离 0.5 差 0.2，离 1.0 差 0.3，舍入到 0.5。量化误差 0.2，相对误差 0.2/0.7 ≈ 28.6%。
图 4：E2M1 数轴：16 个可表示值；0.7 舍入到最近档位 0.5，误差 -0.2（相对误差约 29%）。
单个权重差了四分之一，听着离谱。模型为什么还能用？
直觉的回答是「误差会互相抵消」。这句话值得较真一次。
设一个点积里有 100 万个权重，每个的真实误差要么 +0.2 要么 -0.2（简化），两个方向各一半概率。这些误差加起来，平均会剩多少？
答案是随机游走：剩余误差的典型大小 ≈ √1000000 × 0.2 ≈ 200。而 100 万个权重本身的总量级是百万级——相对误差从单数的 28.6% 被稀释到万分之一量级。
当然，量化误差不是严格随机的（跟数值大小有关），outlier 会系统性拉偏——这正是第 5 章块缩放要解决的问题。但「误差按 √N 稀释」的机制真实存在，它就是大模型敢砍位数的底气。
3.3 第二重保险：网络本身是冗余的
统计稀释之外还有第二重保险：神经网络本身是冗余结构，大量权重的贡献互相重叠。深度学习对单个数字的精度不敏感，对数字之间的相对关系敏感。
两样加起来，才敢说「砍位数可行」。
还有一个值得说的细节：取整的方向。「四舍五入到最近」看似中性，其实有系统性偏差——大量卡在 0.5 边界的值都往上舍，百万参数累积起来，模型输出的分布会被整体推偏。
所以 NVFP4 训练方案里用随机取整（stochastic rounding）：一个值落在两个档位中间时，按距离比例随机取上档或下档；单次看是随机的，期望上舍入误差恰好归零。
这个技巧不占一个比特，却能把长训练里的 loss 曲线压稳。量化工程里这类「看不见但真金白银」的细节，不在公式里，在踩过的坑里。
但 28.6% 同时说明一件事：量化误差的下限由刻度决定，刻度越稀，误差越大。E2M1 只有 8 个正刻度，天生是所有浮点格式里误差最大的。
所以 E2M1 必须配合某种手段，把每个数「落到最近刻度的距离」变相压回来——这个手段就是缩放，第 5 章细讲。在那之前，先插一段：为什么值得这么费劲。
4. 为什么省显存就快：带宽 vs 算力
这一章最容易被「感觉」带偏，所以全用真数字。它回答一个基础问题：省显存为什么能让推理变快。
4.1 先搭个厂房
把显卡想成一个工厂，三样东西：显存是仓库，7B 模型的 FP16 权重 14 GB 货物全堆在这；计算单元是车间，真正干活（乘加）的地方。
仓库和车间之间有一条传送带，宽度固定，行话叫内存带宽，单位是「每秒能搬多少 GB」。
车间产能再高，传送带多宽就是多快——货不过传送带，车间就得空转等。
「车间等货」的状态，行话叫带宽受限（memory-bound）；反过来，货管够、车间排满，叫算力受限（compute-bound）。下面各算一笔账。
4.2 账一：decode 为什么卡在传送带
为什么每个 token 都要把权重搬一遍？这是模型结构决定的：所有权重存在显存里，计算一个 token 时，每一层做的都是「把这个 token 的向量 × 这一层的权重矩阵」。
权重矩阵动辄几百 MB 到几 GB，计算单元片上的缓存只有几十 MB 量级，装不下，只能每处理一个 token 就从显存仓库逐层取货。一次用完，下一个 token 再取一遍——这就是 decode 带宽账的本源。
拿 H200 当例子（141 GB HBM3e 显存，带宽 4.8 TB/s，BF16 算力约 989 TFLOPS，都是官方规格），模型 7B（70 亿参数），FP16 权重 14 GB。
生成一个 token 要做多少活？经验上每个权重大约贡献 2 次乘加：7e9 × 2 = 14 GFLOP。
算力能跑多快？989e12 FLOP/s ÷ 14e9 FLOP/token ≈ 70,000 tokens/s——车间理论满负荷的速度。
带宽能供多快？每生成一个 token，14 GB 权重基本要过一遍传送带：4800 GB/s ÷ 14 GB/token ≈ 342 tokens/s。
两个数字差了 200 倍。车间每秒能吃 7 万个 token 的活，传送带每秒只供得进 340 个——decode 的真实天花板是 342 tokens/s，瓶颈在传送带，不在车间。
图 5：生成一个 token 的两笔账：搬货 14 GB ÷ 4.8 TB/s ≈ 3 ms/token，算货 2×7 GFLOPs ÷ 4000 TFLOPS ≈ 微秒级——搬货比算货慢 1000 倍，瓶颈在搬。
把权重从 FP16 砍到 4-bit（摊上 scale 约 0.56 字节/参数），仓库里 14 GB 变 3.9 GB，传送带同一时刻能多推 3.5 倍的货，decode 上限水涨船高。
这就是「省显存 = 快」的全部机制：快不是因为算得省，是因为搬得少。
（注：这还没算 KV cache。对话越长，每层缓存的 key/value 向量越大，长上下文场景下 KV cache 常占显存的三分之一以上——它同样吃带宽，同样是量化和显存优化的主战场。本文主线是权重，KV cache 不展开。）
4.3 账二：prefill 为什么卡在车间
同样是这个模型，工况变了：你一次性丢进去 4000 个 token 的长文，模型要处理这 4000 个 token 之间的注意力。这时每个权重被复用 4000 次——货搬一次，车间干 4000 份的活。
搬货需求：14 GB，一次，传送带 3 毫秒搬完。干活需求：14 GFLOP × 4000 = 56 TFLOP，车间要干 56 毫秒。
车间的活比货多 18 倍——这次卡脖子的是车间，带宽反而富余。这就是算力受限。
这种工况下，权重砍到 4-bit 省不出多少时间（搬货本来就不是瓶颈）。能救场的是车间本身吞吐翻倍：tensor core 吃 4-bit 数据时乘加速率翻 4 倍，prefill 就直接快 4 倍。
4.4 两条收益线
把两笔账合起来，低精度格式的价值分成互不重叠的两条线。
第一条，省搬运：4-bit 权重让仓库缩 4 倍、传送带需求降 4 倍 → decode 提速、大模型装进小卡。这条线不挑硬件，老卡也吃得到。
第二条，提车间：tensor core 原生吃 4-bit、乘加速率翻倍 → prefill 和训练提速。这条线挑硬件，只有带 4-bit tensor core 的卡才吃得到。
后文所有「值不值」的判断，都可以回到这两条线上对号入座。第 9 章的软件路径和第 10 章的老卡，用的就是这把尺子。
5. NVFP4 格式：两级缩放
这一章讲格式本身：一个 s 管全量会被 outlier 绑架，分块解决「谁绑架谁」，E4M3 还是 E8M0 决定「尺子切多细」，末尾还得再加一级全局 scale。结尾拿真数字把量化全流程走一遍。
5.1 一个 s 管全量：被 outlier 绑架
3.1 的两头崩，解法只有一个方向：缩放。先把数据乘一个比例 s，把整体塞进 0.5~6 的窗口里量化，用的时候乘回去：
真实值 x = 4-bit 档位 q × 缩放系数 s新问题是：一个 s 管多少人？这是所有低精度格式的分水岭。
最省事的方案：整个张量共用一个 s。s 的取法被最大值卡死——为了不溢出，s 必须满足「全张量最大绝对值 / s ≤ 6」。
算笔具体的：16 个权重里 15 个都在 ±0.1，偏偏有一个 30 的 outlier。
s = 30 / 6 = 5（全局比例尺）
15 个小权重：0.1 / 5 = 0.02 → E2M1 最小非零刻度是 0.5 → 全部四舍五入成 015 个有用的权重瞬间清零。姚明进了班级，全班按姚明的尺子量身高——一个极端个体绑架了所有人的精度。
真实模型里 outlier 是常客，单 s 方案在 4-bit 下基本等于自杀。
5.2 分块：16 个人一个 s
把张量切成小块，每块按自己块内的最大值（amax）单独定 s。还是上面那 16 个数，但 outlier 单独占一块：
outlier 块：amax = 30，s = 30/6 = 5，30/5 = 6 → 精确落在档位 6 上
小权重块：amax = 0.1，s = 0.1/6 ≈ 0.0167，0.1/0.0167 = 6 → 也是精确 6两块各用各的尺子，谁也绑架不了谁。块越小，一个 outlier 能绑架的人数越少。
NVFP4 选的块大小是 16 个元素（MXFP4 是 32，这 16 为什么关键，第 7 章细说）。
5.3 块 scale 自己用什么格式存？E4M3 对 E8M0
块 scale 本身也是数，也得占比特。两个候选。
E8M0：8 个比特全是指数，没有尾数。scale = 2^(e-127)，只能取 2 的幂：…、1、2、4、8、16…。256 档，但全是整数倍。
E4M3（FP8 的一种）：1 位符号 + 4 位指数 + 3 位尾数，最大 448。和 E2M1 一样是真正的浮点，能表示 2 的非整数次幂。
图 6：块 scale 的两种「尺子」：E8M0 只能取 2 的幂（间隔全是 ×2），E4M3 刻度密、含 4.5、3 这类非 2 的幂；块内幅值差异大时 E4M3 更精细。
差别在哪，拿块 amax = 2.3 算。scale 要尽量「贴住」amax（贴得越紧，块内动态范围用得越满）：
E8M0：能取的只有 2 和 4
取 2 → 2.3/2 = 1.15，量程利用率 1.15/6 ≈ 19%，浪费
取 4 → 2.3/4 = 0.575，更浪费
E4M3：2.3 = 2¹ × 1.15 → 取 2.25 或 2.375，直接贴住4.5 这个数同理：E4M3 里 4.5 = 2² × 1.125，尾数 001，精确表示；E8M0 里它不存在，只能在 4 和 8 之间二选一，误差 12.5% 起步。
英伟达官方博客给过同数据下两种 scale 格式的量化误差（MSE）对比，E4M3 明显更低——块 scale 用 E4M3，是 NVFP4 和 MXFP4 最本质的一处分歧。
5.4 为什么是两级：E4M3 量程也兜不住全张量
分块解决了「谁绑架谁」，但 E4M3 自己量程有限（最大 448，最小非零 2⁻⁶）。
如果整个张量的 amax 是 1000，s 要算到 1000/6 ≈ 167——还能塞进 E4M3；但训练里激活值的 amax 会漂，有的张量要 s ≈ 4000，E4M3 直接溢出。
NVFP4 的解法是加第二级：整个张量再共享一个 FP32 全局 scale。FP32 量程 10³⁸，大尺度它兜底；E4M3 块 scale 只负责在自己量程内贴住块内 amax。完整重建公式：
x = q × s_block × s_global
图 7：NVFP4 两级缩放：每 16 元素块共享一个 8-bit E4M3 块 scale，全张量再挂一个 32-bit FP32 全局 scale。
量化时两级 scale 怎么算（Transformer Engine 文档公式）：
s_global = 全张量 amax / (448 × 6) ← 448 是 E4M3 最大值，6 是 E2M1 最大值
s_block = (块 amax / 6) / s_global ← 结果转成 E4M3 存（取最近可表示值）读法：全局 scale 先把整个张量的动态范围「对齐」到 E4M3×E2M1 的组合量程，块 scale 再按每块自己的 amax 做精细缩放。两级各管一摊，一个管「大」，一个管「细」。
5.5 开销账 + 和 MXFP4 的完整对比
16 个 4-bit 值（64 bit）摊 1 个 8-bit 块 scale：4.5 bit/元素，外加每张量一个 FP32（可忽略）。比 FP16 省 3.5 倍，比 FP8 省 1.8 倍。
MXFP4 是 32 元素一块、E8M0 scale：4.25 bit/元素，更省 0.25 bit。
| NVFP4（英伟达） | MXFP4（OCP 开放标准） | |
|---|---|---|
| 元素格式 | E2M1 | E2M1（完全相同） |
| 块大小 | 16 | 32 |
| 块 scale 格式 | FP8 E4M3，非 2 的幂，贴得住 amax | E8M0，纯 2 的幂，256 档 |
| 层级 | 两级（块 E4M3 + 全局 FP32） | 一级（就块 scale） |
| 开销 | 4.5 bit/元素 | 4.25 bit/元素 |
| 阵营 | 英伟达自家，CUDA/TE/TRT/vLLM/SGLang 全支持 | AMD/Intel/英伟达共同支持的开放标准，可移植 |
省 0.25 bit 值多少精度？英伟达论文（arXiv:2509.25149）同条件预训练对比：8B 模型下，NVFP4 的 loss 比 16-bit 基线差约 1.5%，MXFP4 差约 2.5%。MXFP4 要多训 36% 的 token 才追平 NVFP4。
0.25 bit 换 1 个百分点，多数场景不划算——这就是英伟达全线押 NVFP4、而 MXFP4 走跨厂商可移植路线的由来。
5.6 全流程走一遍：拿真数字跑
拿真数字把 5.4 的公式完整跑一遍，不跳步。设某张量全局 amax = 12.3，其中一块的块 amax = 2.4：
图 8：拿真数字（全局 amax=12.3、块 amax=2.4）把 NVFP4 量化四步走完：① s_global≈0.00458（FP32 无损）② s_block≈87.4 → E4M3 档位 88（约 0.5% 舍入偏差）③ 元素 ÷0.403，5.96 落到档位 6 ④ 存储 = 16×4-bit + 8-bit 的 88 + 全张量共享的 FP32，4.5 bit/元素。
图 9：NVFP4 离线量化流程：原始权重 → 全张量 amax → s_global（FP32）→ 逐块 amax → s_block（E4M3）→ 逐元素 E2M1 舍入，最终存储布局是 4-bit 数据 + 8-bit 块 scale + 32-bit 全局 scale 三色块。
计算时按 q × s_block × s_global 反向还原。跑一遍你会发现：NVFP4 的量化本身不是黑魔法，就是四步算术。
解压那边同样值得跑一遍：
图 10：同一个 scale（0.403 = 88 × 0.00458）的往返：去程 2.4 ÷ 0.403 = 5.96 → 四舍五入到档位 q=6（压缩，离线，跑一次）；回程 6 × 0.403 还原出 2.418（解压，每个 token 都要走）。误差 +0.018（约 0.75%），来自 87.4→88 和 5.96→6 两处舍入。
黑魔法在下一章——硬件怎么把第 4 步还原里的「×s_block×s_global」变得近乎免费。
6. 硬件到底在加速哪一步：关键路径
这一章回答一个尖锐的问题：NVFP4 的「硬件加速」，加速的到底是哪一步？答案反直觉。
很多人的第一反应：加速算 scale 那套公式吧？
不是。5.4 那两条公式是量化阶段离线跑一次的事：整个张量扫一遍取 amax，公式代进去，s_global 和所有 s_block 算完落盘，结束。一次性的，成本可以忽略。
真正在关键路径上的是每一次矩阵乘（GEMM）时的反量化。回看第 4 章：每生成一个 token 要过几十层，每层几个 GEMM，每个 GEMM 里的每个 4-bit 权重都得先「×s_block×s_global」还原成高精度数，才能参与乘加。
这个动作每个 token 要执行数十亿次，全在出字的关键路径上，一秒都拖不得。
图 11：时间轴上的两条轨道：上轨离线量化只跑一次（成本可忽略），下轨每个 token 的推理要过几十层、每层几个 GEMM，每个 4-bit 权重的反量化全在出字的关键路径上——硬件要加速的就是它。
Blackwell 的做法是把这一步整个融进矩阵乘指令：packed 的 4-bit 数据（显存里两个 4-bit 值挤在一个字节）原封不动喂进 tensor core，拆包、缩放、乘加在流水线内部一气呵成。
中间不存在「还原好的 FP16 张量」这种需要额外读写的中间产物。
图 12：同一个 GEMM 的两条流水线：朴素做法先反量化出 FP16 中间张量（显存/带宽翻倍、多一次访存）再乘加；Blackwell 的做法把 packed 4-bit 原封不动喂进一条 MMA 指令，拆包、乘 scale、乘加在流水线内部一气呵成。
看懂这一步，得先弄清 tensor core 是什么。矩阵乘的原子操作是点积：两列数对应相乘，再求和。
先拿 3 个数手算一遍：
图 13：左边是 3 个数手算 a·b=16（对应相乘、再求和）；右边是矩阵乘的铺法——每个结果格 = A 的一行 · B 的一列，每个格就是一次点积。
矩阵乘就是把点积铺成一片：左矩阵的每一行，对右矩阵的每一列，各做一次点积，铺满结果矩阵。大模型里这两块矩阵动辄 4096×4096，算一笔量级账：
4096×4096 的方阵乘：
结果有 4096×4096 个点积，每个点积 4096 项
总乘加次数 = 4096 × 4096 × 4096 × 2 ≈ 137.4 GFLOP这是一层里一个矩阵乘的活。7B 模型几十层、每层好几个这样的矩阵乘，每生成一个 token 就是几十 GFLOP——第 4 章那个「14 GFLOP/token」就是这么来的。
没有 tensor core 的 GPU 算点积的姿势是：把 a 和 b 搬进寄存器，循环里一次取两个数、乘一次、累加一次，4096 项就是 4096 次「取数-乘-加」串行循环。
算一笔指令账：一个 4096 项点积，CUDA core 要发射 8192 条标量运算指令（4096 乘 + 4096 加），每条还要走取指令、取数、执行、写回的全套流程。同一个点积在 tensor core 里呢——一条 MMA 指令。
8192 比 1，这就是「通用单元是瑞士军刀、tensor core 是装配线」的全部含义：慢不在不会算，在每一步都要花指令去伺候。
tensor core 就是为「把那个循环焊死在电路里」而生的：一条指令端上一小块矩阵（比如 16×8×16 的一片），电路内部自动完成「对应相乘、逐层求和、累加」的整个流水，输出一小块结果。
代价是只能按它规定的姿势喂数据：数据要排成指定形状（fragment）、放在指定位置、走指定指令。
举个具体的：Ampere 时代的 FP16 矩阵乘指令 mma.sync m16n8k16，单条指令完成一个 16×8 输出块的计算：A 侧 16 行、B 侧 8 列、K 方向各 16 个元素。
16×8 = 128 个结果值拆给一个 warp（固定 32 线程的线程组）里的 32 个线程，每线程拿 4 个。数据怎么分给线程、结果落在哪个寄存器，全部写死在指令语义里。
软件 kernel 的大量工作，就是想办法把数据摆成指令想看到的这个姿势。
英伟达从 2017 年 Volta 架构开始每代都在这条装配线上加码：Volta 会 FP16，Turing/Ampere 加了 INT8/INT4，Hopper 加了 FP8，Blackwell 这一代——开始吃 4-bit 浮点，还顺手把「块 scale」这个 NVFP4 的核心机制做进了流水线内部。
这就是本文的主场。
写到这里，一个尖锐的疑问必然冒出来。它问得非常对：
「你说流水线里要乘 scale。那 4-bit 值乘上 scale 之后——比如 s=2——它就不在 0~6 里了，变成 2、3、4、6、12 这种 FP8 甚至更大量级的数。那乘法单元不就得按更宽的格式来算？4-bit 的吞吐优势还有吗？」
这个问题正好戳中 NVFP4 硬件设计里最漂亮的一步棋。答案一句话：硬件从头到尾，没有把 scale 乘到任何单个元素上。
7. 核心：scale 在求和号外面
这一章是全文的心脏：一步中学代数，把第 6 章的矛盾解掉；顺带把 W4A4 双 scale、块 16 和归约树一并讲完。
7.1 一步代数
回看点积：对应相乘，再求和。回看 5.4：还原公式 x = q × s × s_global。
假设硬件「老老实实」干：先把每个元素还原（scale 乘进每个 4-bit 值），再喂乘法单元。结果就是第 6 章那个矛盾——还原后的数是 FP8/FP16 量级，乘法单元必须按宽操作数运转，4-bit 吞吐优势当场清零。
出路藏在一个中学代数事实里：点积是线性的，常数可以提出求和号。
块内所有元素共用同一个 s（这是「分块缩放」的定义），所以对一个块内的点积：
s·q₁×p₁ + s·q₂×p₂ + s·q₃×p₃ + s·q₄×p₄
= s × (q₁p₁ + q₂p₂ + q₃p₃ + q₄p₄)
= s × (原始 4-bit 值的点积)s 在求和号里是常数，提出来。乘法单元里碰到的，永远是原始 4-bit 值；scale 从「每个元素乘一次」变成「每个块乘一次」。
7.2 手算验证
设块内 q = [1, 1.5, 0.5, 2]，p = [2, 1, 1.5, 0.5]，块 scale s = 3（s 取 3 是因为 3 正好是 E2M1 档位；真实场景里 s 是 E4M3 值，道理一样）。
数字全取自 E2M1 真实档位，可按计算器验。
笨办法（先还原再算）：
q 逐个 ×3 → [3, 4.5, 1.5, 6] ← 注意：3、4.5 已不是 4-bit 档位了
3×2 + 4.5×1 + 1.5×1.5 + 6×0.5
= 6 + 4.5 + 2.25 + 3
= 15.75聪明办法（先算后缩放）：
原始值直接点积：
1×2 + 1.5×1 + 0.5×1.5 + 2×0.5
= 2 + 1.5 + 0.75 + 1
= 5.25
再乘一次 s：5.25 × 3 = 15.75 ✓两条路结果一模一样——数学上必然（线性），浮点下也几乎无损（块 scale 只有 8 个比特，乘在部分和上，舍入误差远小于逐元素乘）。
严格说，浮点下两条路并不 bit 级相等：聪明办法先加后乘，笨办法先乘后加，浮点加法不满足结合律，最后一两位会有差异。
但这个差异（相对误差 10⁻⁷ 量级）比量化本身的误差（10⁻² 量级）小三个数量级，工程上可以直接无视。
差别在硬件成本：聪明办法里乘法单元只见过 4-bit 值，s 在 16 个乘加汇完之后才乘一次。
7.3 硬件的四步流水
于是 Blackwell 上这条矩阵乘指令的内部流程变成四步：
1. 窄乘法 lane 执行原始 4-bit × 4-bit 乘加（scaled 值没进过 lane）
2. 16 个元素块的部分和在归约树里汇出来
3. 部分和 × 块 scale（E4M3，一拍），加进主累加器
4. 收尾（epilogue）再乘一次 FP32 全局 scale
图 14：硬件路径 scale 在求和号外面（块部分和只乘一次 scale），软件路径只能逐元素乘 scale——4 倍吞吐差别的根子。
第 4 步多说一句：全局 scale 是整个张量一个常数，对整个输出矩阵一视同仁，所以它不用挤进流水线，在结果写回前的收尾阶段乘掉就行。每块输出只摊一次，成本约等于零。
7.4 W4A4：两侧都有 scale
权重和激活都是 4-bit 的工况（行话 W4A4）也照样成立：激活侧有自己的块 scale s_a，但两侧沿同一个方向（K 维）分块、块大小一致。
s_a 和 s_w 在块内都是常数，一起提出求和号，块点积结果乘 s_a × s_w 即可。
图 15：W4A4 双 scale 对齐：A 的一行 16 元素与 W 的一列 16 元素对齐成同一个 16×16 小块，块部分和 × (s_a·s_w)——两侧 scale 都是块内常数，一起提出求和号。
数学结构完全不变——「常数提出求和号」这件事，不关心常数有几个。
7.5 块为什么是 16（推断）
推断（英伟达未公开 tensor core 微架构，以下为基于公开资料的反推）：块大小 16 很可能不纯粹是精度选择，还与乘加阵列内部归约树的粒度对齐——scale 乘法必须插在「块边界」上，归约树每一级的汇合点决定了 scale 能插刀的位置，16 恰好让每个块的部分和独立成段、各自乘一次 scale。「scale 提出求和」本身是由 NVFP4 两级缩放定义 + 4 倍吞吐数据推出的数学必然性；「具体在流水线哪一级乘」属于推测。MXFP4 的块 32 同理对齐其硬件粒度。
这一步棋同时解答了第 6 章那个疑问：正因为 scale 被挪到了求和号外面，乘法 lane 里流转的永远是 16 个档位之内的 4-bit 值——「4-bit 进 lane」和「要乘 scale」两件事不再打架。
8. 4 倍吞吐从哪来：lane 切分
这一章解释最后一块：同一块硅，喂 4-bit 时吞吐为什么翻 4 倍。
「FP4 tensor core」这个说法很容易让人以为：硬件在用 4-bit 精度做数学。
不是。和 INT8/FP8/FP16 张量核心一样，FP4 tensor core 的累加器仍然是 FP32——最终结果精度一点没缩水。
「FP4」指的是输入格式和吞吐档位：指令接受 packed 4-bit 操作数，每个周期吞掉的元素数是 BF16 档的 4 倍。
先立一个对照：B200 官方标称的 dense 吞吐——
FP4 : FP8 : BF16 = 9 : 4.5 : 2.25 PFLOPS正好 4 : 2 : 1。注意 FP8 对 BF16 也是 2 倍——「位数减半、吞吐翻倍」这条规律贯穿整条低精度路线。它说明这不是 FP4 的个别魔法，是 tensor core 的通用构造方式。
推断（英伟达未公开 lane 级微架构，以下为最合理的解释）：机制是乘法 lane 的切分。拿高速公路类比展开说：一条 8 车道的高速公路，每车道跑的是 FP16×FP16 乘法器——造一条这样的车道贵（16 位×16 位的部分积，电路规模不小）。
现在把同一条路改造成 32 条窄巷，每条巷跑 4×4-bit 乘法器——4 位×4 位的部分积只有 16×16 的约 1/16 规模，同一片地皮、同一份功耗预算，塞得下的车道数翻了 4 倍。
每车道每周期仍然完成一个元素的乘加，车道数 ×4，总吞吐 ×4。
图 16：同一块硅的两种 lane 切法：4N 条窄 lane（4×4 乘法器）替代 N 条宽 lane（16×16 乘法器），同样面积功耗、吞吐 ×4。
换句话说：不是「4-bit 数学更便宜」，是「同一块硅按 4-bit 切能切出 4 倍数量的乘法单元」。
FP8 的 2 倍同理（8 位乘法器约 1/4 成本，塞 2 倍车道）。
先把 9 PFLOPS 这个数的口径说清楚：它是 dense（无稀疏）标称值。官方营销里也常用 sparse（结构化稀疏，理论再翻倍）的数字，看参数时要对齐口径。本文全部用 dense，对比只在同口径下成立。
这个切法还解释了为什么累加器必须是 FP32：窄 lane 算出的部分积精度低，靠高精度累加器把精度补回来——「低精度输入 × 高精度累加」是 tensor core 一贯的配方，FP4 只是把输入端切到了最细。
到这里，前面各章的悬念全部闭环：缩放不是加速对象，每次 GEMM 的反量化才是；反量化能融进 4-bit 流水线，是因为 scale 被提到了求和号外；4 倍吞吐成立，是因为 lane 按 4-bit 切分，而 lane 里流转的永远是原始 4-bit 值。
9. 对手视角：软件路径三档
这一章看软件侧。说「软件写 kernel」，其实有三条路，吞吐上限完全不同：
①② 两档都是通用单元反量化、喂 16-bit tensor core 收尾，封顶原生速率的 1/4；③ 档用 inline PTX 直驱 4-bit 矩阵乘指令，能摸到原生速率，但工程成本是 CUTLASS 级别。
档与档的差别不在优化功力，在能摸到哪个天花板。
图 17：软件路径三档：① 朴素反量化 ② 延迟缩放（hoisted scaling）③ 原生指令直驱（inline PTX tcgen05），吞吐上限分别是 FP16 速率、FP16 速率、FP4 速率。
9.1 第一档：朴素反量化
kernel 里把 E2M1 查表转成 FP16（16 个值，查表而已），逐个元素乘上 s_block×s_global，还原成 FP16 张量，喂给 FP16 tensor core 做常规矩阵乘。
老显卡（A100/H100）跑 4-bit/INT4 量化模型，标准姿势就是这个。
精度完全正确，但每个元素多一次通用单元的乘法，吞吐封顶 16-bit lane 速率——第 8 章那 2.25 PFLOPS 一档。
9.2 第二档：延迟缩放——把第 7 章的数学学过来
关键观察：E2M1 的 16 个值转成 FP16 是严格无损的（1.5 就是 1.5，不存在表示误差）。
所以软件可以完全复刻硬件的手法：把原始值（不带 scale）转 FP16 喂 MMA，让每条 MMA 指令恰好覆盖一个 16 元素块、产出一个块的部分和，再在向量单元里乘一次块 scale、FMA 进累加器。
有个对得上的细节：Ampere 上 FP16 的 mma.sync m16n8k16 指令，单条指令正好处理 K 方向 16 个元素——和 NVFP4 块大小天然吻合，kernel 写成「一条 MMA → 乘 scale → FMA」的循环就很干净。
这一档精度上和硬件路径数学等价（同一个 7.1 的恒等式）。
代价：吞吐依然封顶 16-bit lane 速率（约原生的 1/4），外加每个块打断一次累加器的 rescale 开销。
省掉的是最便宜的部分（每 16 个元素一次的 scale 乘），变不出来的是最贵的部分（4 倍 lane 切分）。
9.3 第三档：原生指令直驱
用 inline PTX 直接发射 tensor core 的 4-bit 矩阵乘指令（Blackwell 上叫 tcgen05），让硬件自己走 7.3 那套流水线。
这一档就是原生路径本身——cuBLASLt 和 CUTLASS 底层干的就是这件事。
全部 4 倍吞吐拿到手，但门票很贵：packed 数据要按指定布局放进指定存储（tensor memory/共享内存）、scale 张量要按规矩摆放、warp 级调度有一堆约束。
CUTLASS 级别的工程，几乎没有团队手写裸指令，都是调现成库。
9.4 为什么 ①② 封顶 1/4
因果链一句话：自由 CUDA 代码摸不到 4-bit 乘法 lane。
4-bit 乘法器只对专用矩阵乘指令开门，通用计算管线里压根没有「4-bit 浮点乘法」这个数据类型。
你的 4-bit 值进乘法单元之前必须先变成 16-bit 形态，而 16-bit lane 的上限就是原生 4-bit lane 的 1/4。
软件能学走数学，学不走电路。
9.5 一个现实注脚
目前落地最多的其实不是 W4A4，而是 W4A16（权重 4-bit、激活 16-bit）。
原因在第 4 章已经埋了——权重是静态的，amax 提前算好存着就行；激活跟着输入实时变化，分布没法提前知道，只能推理时现估现量化，代价和风险都高一截。
所以工程上的顺序是：先把权重砍到 4-bit 吃搬运收益（decode 提速立竿见影），W4A4 等基础设施成熟再上。
这也解释了为什么 RTX 50 的 tensor core 有 4-bit 能力，生态却主推 weight-only。
10. 没有 Blackwell 的卡，4-bit 还有没有价值
这一章看老卡：Hopper 和 Ampere 没有浮点 4-bit（FP4）指令，但 4-bit 不是死路。
先泼冷水：PTX 是可移植汇编层，上限就是目标架构自己的指令集。
Hopper 能摸到 FP16/BF16/FP8/INT8 的 wgmma 指令，Ampere 是 mma.sync，浮点 4-bit（FP4）专用矩阵乘指令在这些架构上物理上不存在（整数 INT4 指令 Turing 起就有，第 6 章提过，但 FP4 浮点值没法直接喂给它——W4A16 还是「反量化 + 本代乘加」）——第 8 章那 4 倍变不出来。
但 4-bit 的价值是两条线（4.4），老卡能吃掉其中一条半。
图 18：老卡（Hopper/Ampere）4-bit 路径（没有浮点 4-bit FP4 指令，整数 INT4 从 Turing 起就有）：packed 4-bit 存储（显存收益）→ kernel 内 E2M1→FP16 无损查表 → 块 scale → 本代 tensor core（wgmma/mma.sync）。batch=1 decode 是带宽瓶颈：权重搬运 14 GB → 3.9 GB，比值 3.6 倍与卡无关（342 → 约 1200 tokens/s 是文章 H200 4.8 TB/s 的账）；A100 端到端实测 1.8~2.9 倍（vs FP16，Marlin 论文 Table 2）；prefill / 大 batch decode 是算力瓶颈，封顶本代 16-bit 速率。
第一条，省搬运，整条拿到：E2M1 packed + E4M3 块 scale 就是普通显存数据，任何架构都能读。
仓库缩 4 倍、传送带需求降 4 倍，不挑硬件。decode 场景直接受益——4.2 那笔账里，14 GB 变 3.9 GB，342 tokens/s 的上限跟着往上走。
这不是纸面账：A100 上 vLLM 端到端实测，4-bit 权重推理比 FP16 快 1.8~2.9 倍（Marlin 论文 Table 2：Yi-34B 2.90x、Llama-2-70B 四卡 2.02x、Falcon-180B 八卡 1.76x）；kernel 级能摸到 3.6 倍上限附近，端到端被 KV cache、attention、norm 这些没量化的部分拖下来。
第二条，提车间，封顶本代：kernel 内用第 9 章档②的延迟缩放手法把反量化开销压到最低，喂本代 tensor core，榨干 16-bit/8-bit 速率。
prefill/训练不会比非量化更快，但显存省 4 倍。
补充一条工程细节：Ampere 老卡上很多 INT4 kernel（包括 Marlin）的乘加并不走 16-bit 浮点 lane，而是把 4-bit 值映射成 INT8、走 INT8 tensor core——INT8 在 Ampere 上是 2 倍于 FP16 的档位，反量化后的乘加也能多赚 2 倍。
同一张卡，选对档位就是 2 倍差距。
介于两条线之间：老卡上 4-bit 最大的实际收益其实是容量。
拿消费级算笔账：RTX 4090 显存 24 GB，跑 FP16 的 24B 模型需要 48 GB，装不下；4-bit（4.5 bit/参数）只要约 13.5 GB，稳稳装下还能留 KV cache 的地方。
同一张卡，从「最多 12B」变成「24B 随便放」，这是「装得下」的胜利，不是「跑得快」的胜利——但对很多人来说，装得下比跑得快重要得多。
数据中心侧同理：单张 80 GB 的 H100，FP16 下勉强装下 30B 模型（权重约 60 GB，剩 20 GB 给激活和 KV cache），70B 完全没戏。
换 NVFP4，70B 权重压到约 39 GB，一张卡装下还有富余。
很多真实部署里，「单卡跑 70B」不是性能优化，是能不能跑的区别。
但代价是有的，不记账的话价值判断不成立。
单个 4-bit 权重的误差最大能有 28.6%（第 3 章的 0.7→0.5）。它炸不塌整个网络，机制和 5.6 一样：误差在权重之间近似不相关、有正有负，一次点积求和几万项，互相抵消。最终答案的误差远小于单个权重的误差。
实测损失也不大（全部是 4-bit 权重 + 16-bit 激活，对比 FP16）：
大模型的困惑度代价极小：OPT-175B 只升 0.03（8.34→8.37）、BLOOM-176B 升 0.10（8.11→8.21）、Llama-2-7B 升 0.13（5.47→5.60，AWQ）（GPTQ/AWQ 论文）。真实生态里在跑的 llama.cpp Q4_K_M，Llama-2-7B 困惑度升 0.08 点（+1.4%）、Llama-3-8B 升 0.15~0.18 点（+2.4%~2.8%）（llama.cpp 官方 scoreboard）。
下游任务掉几分：LLaMA-7B 的 MMLU 38.41→37.71（AWQ 4-bit）或 35.39（GPTQ 4-bit）；Marlin 论文的 Llama-2-7B INT4，MMLU/WinoGrande/ARC 三项均值 56.96→53.63（−3.33）。
也就是说，7B~175B 这个量级，4-bit 的精度代价是「困惑度升一点、benchmark 掉 1~3 分」——对比 4 倍显存和 2~3 倍速度，是小损失。
代价陡增的情况有两种：位宽降到 3-bit，或者模型变小。
3-bit 的损失陡了：OPT-175B 困惑度升 0.34（8.34→8.68）、LLaMA-7B 的 MMLU 掉 3~8 分。而且不做 GPTQ/AWQ 的误差补偿、只朴素舍入（RTN）的话，OPT-175B 困惑度直接从 8.34 爆到 7300——模型直接废了。这正是 GPTQ/AWQ 存在的原因。
小模型吸收误差的冗余少：OPT-1.3B 4-bit 困惑度升 0.84，是 175B 那个 0.03 的二十多倍。
所以「4-bit 还有没有价值」的完整答案是：速度和显存是确定收益，精度是小损失，小损失随模型变小、位宽变低而放大。大模型 + 4-bit 及以上，这笔买卖基本不亏；小模型 / 3-bit / 要求苛刻的场景，先测了再说。
这条路上生态非常成熟：A100 上的 Marlin（vLLM/AutoGPTQ 体系里长期在用的 INT4 GEMM kernel）、vLLM 和 SGLang 的 W4A16 支持，都是「packed 低精度存储 + 软件反量化 + 本代 tensor core」的不同实现。
消费级的 RTX 50 系列（Blackwell 的 SM120）也带 4-bit tensor core，但完整的 W4A4 需要数据中心级 SM100+，消费卡目前基本走 weight-only 路线。
消费卡上 4-bit 最先落地的其实不是 LLM，是图像生成：2025 年 5 月英伟达用 TensorRT 把 Stable Diffusion 一类模型的部分层转成 FP4 跑在 RTX 50 上，单卡出图速度明显提升。
图像生成模型小、算力受限占比高，正好吃到 4-bit 的乘加吞吐收益，算是消费级验证 4-bit 流水线的先锋场景。
11. NV 的开放边界
最后一章正文：英伟达开放了什么、没开放什么，以及「你」到底能做什么。
一句话概括：门是开着的，门后只有一条固定传送带。
开放的部分：PTX 指令集是公开文档，4-bit 矩阵乘指令（tcgen05 及其 block-scaled 变体）白纸黑字写在里面。
CUTLASS 是开源库，任何人都能照着驱动 4-bit tensor core，甚至能读它怎么摆布局、怎么调度。
不存在「闭源指令集」那堵墙。想研究、想移植，材料都是现成的。
不开放的部分：开放的粒度停在「指令」这一层。
操作数格式、块大小（16 或 32）、累加精度——这是出厂定好的菜单，点菜可以，研发新菜不行。
流水线内部——lane 怎么切、scale 在哪一级乘、归约树多深——是固定电路，不能自定义组合，也没有用户可写的微代码。
你不能「伸手进流水线内部」只借用那 4N 条窄 lane、配自己随意的数据流。
其实这是 tensor core 自 2017 年 Volta 以来的一贯模式：矩阵乘指令从来都只是 PTX 级的黑盒，没人能自由编排它的内部。
设计哲学很直白：固定功能换吞吐。
从商业逻辑看，这也是一道防线：4-bit tensor core 的完整能力锁在 SM100+ 的数据中心卡上，消费级只给一部分，配合 CUDA/CUTLASS 生态的绑定，「想用满 NVFP4 的全部性能」这件事的答案永远是买英伟达最新的数据中心卡。
开放标准 MXFP4 之所以有 AMD、Intel 参与推，也是同一个背景下各家在找自己的破口。
通用 CPU 给你想写什么就写什么的自由，矩阵乘很慢；tensor core 把深度学习最高频的负载（块点积）焊死在电路里，换回成倍的吞吐。
业界想定制时的标准答案因此只有一个：把定制逻辑挪到盒子外面——MMA 之前的数据摆放和预处理、MMA 之后的收尾运算（bias、激活、再缩放随便写）、kernel 结构的调度模板。
盒子保持固定，灵活性来自盒子外围。
落到「我到底能做什么」，三件事：第一，CUTLASS 官方示例里有现成的 4-bit block-scaled GEMM 模板，读它的布局摆放和调度写法，是理解这套硬件最好的一手材料。
第二，有菜单外的定制需求（比如换个 scale 格式），走第 9 章 ①② 档软件路径，拿 1/4 吞吐换完全自由。
第三，量化侧（5.6 那四步离线流程）完全自主，scale 的算法可以按自己的分布特征改；只有消费侧的姿势，是硬件焊死的。
速查表
| 问题 | 答案 |
|---|---|
| E2M1 能表示几个数？ | 16 个：0、±0.5、±1、±1.5、±2、±3、±4、±6，量程 ±6，裸用不可行 |
| NVFP4 是什么？ | E2M1 元素 + 每 16 元素一个 FP8 E4M3 块 scale + 每张量一个 FP32 全局 scale，4.5 bit/元素 |
| NVFP4 和 MXFP4 差在哪？ | 元素数值完全相同；NVFP4 块更小（16 vs 32）、scale 更细（E4M3 非 2 的幂 vs E8M0 纯 2 的幂）、多一级全局 scale；同条件预训练 loss 差 1.5% vs 2.5% |
| 省显存为什么快？ | decode 是带宽受限（7B 例子：算力够 70,000 tok/s，带宽只够 342 tok/s），4-bit 让搬货量降 4 倍 |
| 硬件加速的是哪一步？ | 不是离线算一次 scale，是每次 GEMM 关键路径上的反量化 |
| 4 倍吞吐从哪来？ | lane 切分：同一块硅按 4-bit 切出 4 倍数量的窄乘法 lane（推断）；B200 标称 9:4.5:2.25 PFLOPS |
| 软件能复刻吗？ | 数学能（延迟缩放，精度等价），电路不能（封顶原生 1/4）；拿满速率须直驱原生指令 |
| 老卡能用 4-bit 吗？ | 能：搬运/容量收益整条拿到，算力收益封顶本代 tensor core |
| 4-bit 精度代价多大？ | 小：7B~175B 困惑度 +0.03~0.13、benchmark 掉 1~3 分；模型越小、位宽越低损失越大，3-bit 是陡坎 |
12. 收尾
这篇文章从 16 个数开始，停在一条传送带上。
NVFP4 的价值不在 16 个数本身——谁都有——在两级缩放和硬件流水线的配合。「常数提出求和号」是数学内核，lane 切分是物理内核，两个内核咬合才有 4 倍。
低精度竞赛的评分标准是两条独立的收益线：省搬运 / 提车间。评估任何量化方案之前，先问它在这两条线上各拿了多少。
「软件自定义算子打不过原生指令」，也不是谁不努力——是固定功能加速器的定价逻辑。你一开始就没买自由度。
参考文献
- NVIDIA 博客 · Introducing NVFP4 for Efficient and Accurate Low-Precision Inference（2025-06-24）：E2M1 值集、两级缩放定义、NVFP4/MXFP4/裸 FP4 格式对比、E4M3 vs E8M0 误差示例
https://developer.nvidia.com/blog/introducing-nvfp4-for-efficient-and-accurate-low-precision-inference/ - NVIDIA 博客 · NVFP4 Trains with Precision of 16-Bit and Speed and Efficiency of 4-Bit（2025-08-25）：NVFP4 训练 recipe
https://developer.nvidia.com/blog/nvfp4-trains-with-precision-of-16-bit-and-speed-and-efficiency-of-4-bit/ - NVIDIA 论文 · Pretraining Large Language Models with NVFP4（arXiv:2509.25149）：NVFP4 vs MXFP4 同条件预训练对比（8B 模型、1.5% vs 2.5% loss 差、MXFP4 需多训 36% token）
https://arxiv.org/abs/2509.25149 - NVIDIA 博客 · Blackwell Delivers World-Record DeepSeek-R1 Inference Performance（2025-03-18）：DGX B200 上 FP4 跑 671B 的吞吐数据
https://developer.nvidia.com/blog/nvidia-blackwell-delivers-world-record-deepseek-r1-inference-performance/ - Transformer Engine 官方文档 · NVFP4（两级缩放公式 s_global/s_block、设备支持矩阵）
https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/features/low_precision_training/nvfp4/nvfp4.html - OCP MX 联盟论文 · Microscaling Data Formats for Deep Learning（arXiv:2310.10537）：MXFP4 = E2M1 + E8M0 + 块 32 的规范定义
https://arxiv.org/abs/2310.10537 - NVIDIA PTX ISA 文档：tensor core 矩阵乘指令（含 4-bit / block-scaled 变体）
https://docs.nvidia.com/cuda/parallel-thread-execution/ - CUDA 12.8 Release Notes：cuBLASLt micro-scaled GEMM（E2M1 + UE4M3 scale + 16 元素块 = NVFP4，要求 compute capability 10.0+）
https://docs.nvidia.com/cuda/archive/12.8.0/cuda-toolkit-release-notes/ - vLLM · LLM Compressor 文档 · FP4 Quantization with NVFP4（W4A4 量化、SM100 门槛、老卡 weight-only 回退）
https://docs.vllm.ai/projects/llm-compressor/en/latest/examples/quantization_w4a4_fp4/ - SGLang 文档 · Quantization（NVFP4 在线量化、NVFP4→MXFP4 跨厂商重量化）
https://docs.sglang.ai/advanced_features/quantization.html - NVIDIA 博客 · TensorRT Unlocks FP4 Image Generation for RTX 50 Series（2025-05-14）：消费级 Blackwell 的 4-bit 支持范围
https://developer.nvidia.com/blog/nvidia-tensorrt-unlocks-fp4-image-generation-for-nvidia-blackwell-geforce-rtx-50-series-gpus/ - GPTQ 论文 · GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers（arXiv:2210.17323）：OPT/BLOOM 4-bit/3-bit 困惑度（OPT-175B 4-bit +0.03、3-bit +0.34；RTN 3-bit 崩溃至 7300）
https://arxiv.org/abs/2210.17323 - AWQ 论文 · AWQ: Activation-aware Weight Quantization for On-Device LLM Compression and Acceleration（MLSys 2024 / arXiv:2306.00978）：Llama-2/LLaMA INT4/INT3 困惑度与 MMLU/HellaSwag/ARC 数据
https://arxiv.org/abs/2306.00978 - Marlin 论文（arXiv:2408.11743）：A100 端到端加速 vs FP16（Table 2）、Llama-2-7B INT4 精度（Table 1）
https://arxiv.org/abs/2408.11743 - llama.cpp 官方 perplexity scoreboard（tools/perplexity/README）：Q4_K_M/Q3_K_M/Q2_K 相对 FP16 的困惑度数据
https://github.com/ggml-org/llama.cpp/blob/master/tools/perplexity/README.md