AI 芯片浮点单元（FPU）与低精度数据通路深度解析：从 IEEE 754 舍入与 FMA 运算到 FP8/FP4 块缩放如何决定 Tensor Core 的算力上限核心痛点：算力榜单上有一组每年都在刷新、却很少有人认真追问的数字——B200 的 FP4 稠密算力 7702.5 TFLOPS，FP64 只有 44.8 TFLOPS，两者相差 177 倍；而跑一个 MMA 指令，FP4 用 12.6 个时钟周期，FP64 用 11.2 个，只差 1.27 倍。同一个芯片、同一套 Tensor Core，为什么吞吐能差两个数量级，延迟却几乎一样？更让人费解的是另一组数字：2025 年底的 B300 把数据中心 GPU 保持了十五年的 FP64:FP32 = 1:2 一刀砍到 1:64，峰值 FP64 从 B200 的 37 TFLOPS 掉到 1.2 TFLOPS——相当于把一整块 HPC 能力从旗舰芯片上抹掉。这不是工艺退步，也不是设计失误。真正要回答的问题是：AI 芯片的浮点数据通路到底是怎么被设计出来的？为什么"加宽"比"加深"更划算？为什么累加器精度比乘法精度更卡脖子？以及当 FP8/FP4 加块缩放成为主力之后，新的瓶颈又转移到了哪里？适配人群：做推理/训练性能优化的工程师；写 Triton/CUTLASS 内核、需要理解块缩放 MMA 硬件约束的算子开发者；负责模型量化选型（FP8 还是 NVFP4、块大小取 16 还是 32）的技术决策者；以及对"浮点数在硅片上到底长什么样"感兴趣的底层系统与芯片方向读者收获能力：建立"浮点硬件是一条用精度换吞吐的工程流水线"的心智模型；吃透 IEEE 754 的位域经济学（符号/指数/尾数如何瓜分位宽）与 FMA 单次舍入的硬件价值；理解为什么吞吐扩展靠加宽数据通路而非加深流水线（177x 吞吐 vs 1.27x 延迟的微架构解释）；掌握精度格式完整谱系（FP32/TF32/FP16/BF16/FP8/FP6/FP4）的位域、动态范围与量化误差；吃透 MXFP4 与 NVFP4 两种块缩放方案的硬件差异与误差来源；理解块缩放 MMA 在 TMEM 与缩放因子布局上的真实工程约束；最终获得一份可直接用于项目的精度选型决策表与八条踩坑记录技术背景与演进逻辑：从"精度是标准"到"精度是旋钮"要理解今天的 AI 芯片，先要接受一个反直觉的前提：浮点运算的每一个精度档位，本质上都是芯片设计者用硅面积和功耗买来的一档吞吐——精度不是天经地义的常数，而是一个可调、可砍、可换的工程参数这个认识花了四十年才成为行业共识，中间经历了三个阶段阶段一（1985-2016）：精度是"标准"——IEEE 754-1985 把浮点格式与舍入规则统一下来，此后三十年的主旋律是"算得准"：FMA 在 2008 年被纳入标准并进入硬件，双精度是科学计算不可动摇的底线；这一阶段浮点单元的优化目标是正确性，不是吞吐阶段二（2017-2022）：精度第一次变成"旋钮"——Volta 的第一代 Tensor Core 把 4x4x4 矩阵乘做成硬件指令，Ampere 引入 TF32（19 位）让训练第一次接受"有损但够用"的精度，Hopper 把 FP8 做进 Transformer 引擎；这一阶段的关键转变是：模型开始容忍噪声，硬件因此获得了把精度换成并行度的许可阶段三（2023-2026）：精度变成"体系"——OCP 微缩放（Microscaling）格式把"块共享指数"标准化，Blackwell 第五代 Tensor Core 原生支持 FP4/NVFP4，块缩放 MMA 指令成为训练主力；与此同时，B300 把 FP64 砍到 1:64，宣告当 AI 收入占据绝对主导时，连"精度是标准"这条最后的底线也会为 AI 吞吐让路演进时间线（text 树）：AI 芯片浮点数据通路的四十年 ├── 1985: IEEE 754-1985 发布 │ ├── 贡献: 符号/指数/尾数三段编码 + 四种舍入模式统一, 浮点结果可移植 │ └── 定位: 精度第一次成为"标准", 而不是各家私有的近似 ├── 2008: IEEE 754-2008 纳入 FMA │ ├── 贡献: rn(X×Y+Z) 单次舍入, 成为深度学习最核心的原子操作 │ └── 硬件: NVIDIA GPU 自 compute capability 2.0 起硬件支持, 当年 x86 CPU 尚无 ├── 2010: Fermi 消费卡 FP64:FP32 = 1:8 (硬件 1:2, 驱动限流) │ ├── 背景: 数据中心卡保持 1:2, 精度成为市场分割杠杆 │ └── 结果: 消费卡 FP64 一路走低: Kepler 1:24 - 2014 年 1:32 - Ampere 1:64 ├── 2017: Volta 第一代 Tensor Core │ └── 贡献: 4x4x4 矩阵乘做进硬件, 精度第一次成为可编程的吞吐旋钮 ├── 2020: Ampere 引入 TF32 (19 位); 消费卡 FP64 定格 1:64 │ └── 贡献: 训练第一次接受"低于 FP32"的近似精度 ├── 2022: Hopper FP8 (E4M3 / E5M2) + Transformer 引擎 │ └── 贡献: 8 位正式进入训练与推理主流 ├── 2023: OCP 微缩放格式 (MXFP8 / MXFP6 / MXFP4) │ ├── 贡献: 32 元块共享 E8M0 幂次指数, 4 位在训练中可用 │ └── 代价: 幂次缩放粒度粗, 块内误差偏大 ├── 2024: Blackwell 第五代 Tensor Core: FP4 / NVFP4 原生 │ ├── 贡献: 16 元块 + E4M3 缩放 + 每张量 FP32 二级缩放 │ └── 实测: DeepSeek-R1-0528 上相对 FP8 精度损失 = 1% ├── 2025: cuBLAS 支持 FP64 模拟 (Ozaki); B300 把 FP64 砍到 1:64 │ ├── 数据: 峰值 FP64 从 B200 的 37 TFLOPS 降至 B300 的 1.2 TFLOPS │ └── 信号: AI 收入占比高到足以砍掉传统 HPC 精度 └── 2026: 块缩放 MMA 成为训练主力 (tcgen05.mma.block_scale) ├── 能力: MXFP8/MXFP6/MXFP4/NVFP4 原生, 2~4x BF16 吞吐 └── 现状: MMA 变快后, 量化开销与 SFU 成为新瓶颈值得单独拎出来说的是FP64 的消亡，因为它是"精度作为旋钮"最锋利的证据：消费级 GPU 的 FP64:FP32 从 2010 年 Fermi 的 1:8 一路降到 2020 年 Ampere 的 1:64；从 GTX 480（2010）到 RTX 5090（2025）十五年，消费卡 FP64 只涨了 9.65 倍（0.17 - 1.64 TFLOPS），同期 FP32 涨了 77.63 倍（1.35 - 104.8 TFLOPS）但真正标志性的转折发生在数据中心卡上：数据中心 GPU 一直守着 1:2 或 1:3 的 FP64:FP32 比例，直到 B300（Blackwell Ultra）把它降到1:64——和消费卡同一个比例；这不是技术做不到，而是当绝大多数收入来自 AI 且 AI 训练根本不用 FP64 时，维持高精度单元的动机就消失了替代方案已经就位：既然低精度 Tensor Core 过剩，那就用它们来"拼"出高精度——Dekker（1971）的双-float 方案把 FP64 拆成两个 FP32（A = a_hi + a_lo，有效尾数 48 位而非 53 位），Ozaki 方案更激进：把 FP64 拆成多个 FP8 块，用 Tensor Core 算出所有 A_i×B_j 再在 64 位精度下求和；NVIDIA 已于 2025 年 10 月在 cuBLAS 中加入该支持这条演进线的结论很清晰：AI 芯片的浮点单元已经从"通用数值计算部件"演化成了"按精度分档的吞吐流水线"——每一档精度对应一个硅面积预算、一个数据通路宽度、一个能效点；理解这条流水线，就理解了 AI 算力的上限从何而来本章结论：接下来的六道原理题必须逐个答清楚——位宽怎么分配（位域经济学）、乘加怎么做成一个原子（FMA）、吞吐怎么扩（加宽还是加深）、瓶颈在哪（乘法还是累加）、4 位怎么还能用（块缩放）、以及 FP64 走了之后高精度怎么办（模拟）核心原理深度解析：浮点硬件的六道原理题原理一：位域经济学——符号、指数、尾数如何瓜分位宽IEEE 754 把二进制浮点数编码成三个字段：1 位符号、若干位指数（带偏移 bias）、若干位尾数（significand，且整数部分的 1 隐含不存）；这三段位宽怎么分，直接决定这个格式的"能力形状"（text 树）：浮点位域经济学: 三段位宽决定三种能力 ├── 符号位 (1 bit, 固定) │ └── 决定: 是否有负数 (所有 AI 格式都有) ├── 指数位 (e bits) │ ├── 决定: 动态范围 (能表示的最大值 / 最小值) │ ├── 每多 1 位指数 - 动态范围翻倍 │ └── 代价: 尾数少 1 位, 精度减半 └── 尾数位 (m bits) ├── 决定: 相对精度 (相邻可表示数的间距) ├── 相对误差上界: 2^(-(m+1)) └── 每多 1 位尾数 - 精度翻倍, 但范围减半由此得到 AI 精度格式的完整谱系——注意指数位与尾数位此消彼长的取舍（text 树）：AI 浮点格式谱系 (位域 - 能力) ├── FP32 (1-8-23) │ ├── 动态范围: ~1e-38 ~ 3.4e38 │ ├── 相对精度: 2^-24 ≈ 6e-8 │ └── 定位: 累加与主权重的事实标准 ├── TF32 (1-8-10, Ampere 引入) │ ├── 动态范围: 与 FP32 相同 (指数位不变) │ ├── 相对精度: 2^-11 ≈ 5e-4 (尾数砍到 10 位) │ └── 定位: 训练中"几乎无损"的 19 位近似 ├── FP16 (1-5-10) │ ├── 动态范围: ~6e-5 ~ 65504 (指数只有 5 位, 是主要短板) │ └── 定位: 推理主力; 训练需 loss scaling 防下溢 ├── BF16 (1-8-7) │ ├── 动态范围: 与 FP32 相同 │ ├── 相对精度: 2^-8 ≈ 4e-3 (比 FP16 粗 8 倍) │ └── 定位: 训练首选, 因为动态范围比精度更卡脖子 ├── FP8 E4M3 (1-4-3) / E5M2 (1-5-2) │ ├── E4M3: 范围 ±448, 精度 2^-4; 用于权重与激活 │ ├── E5M2: 范围更大精度更粗; 用于梯度 │ └── 定位: Hopper 起训练与推理的通用低精度 ├── FP6 E3M2 / E2M3 │ └── 定位: FP8 与 FP4 之间的过渡档, 由块缩放 MMA 原生支持 └── FP4 E2M1 (1-2-1) ├── 可表示值: 0, 0.5, 1, 1.5, 2, 3, 4, 6 (及负值), 范围约 ±6 ├── 相对精度: 2^-2 = 0.25 (最粗) └── 定位: 必须配合块缩放才能用一个必须建立的直觉：BF16 的尾数比 FP16 少 3 位（精度差 8 倍），但训练界几乎一致选 BF16——原因就在这张表里：BF16 的 8 位指数与 FP32 完全一致，意味着它不会在梯度下降时突然下溢；动态范围不够会毁掉训练，精度不够只会让训练慢一点，这是 AI 场景对浮点格式的第一条硬偏好原理二：FMA——AI 算力的原子操作，以及它为什么必须是"单次舍入"深度学习最核心的运算是矩阵乘加，拆到最底层就是无数个"乘完再加"；问题是：这个操作到底是两次舍入还是一次？非融合的写法是先把乘积累积到中间结果（舍入一次），再与累加器相加（再舍入一次）：m a t h r m r n ( m a t h r m r n ( X × Y ) + Z ) mathrm{rn}(mathrm{rn}(X × Y) + Z)mathrmrn(mathrmrn(X×Y)+Z)FMA 的写法是把乘积算成双倍宽度后直接与 Z 相加，只舍入一次：m a t h r m F M A ( X , Y , Z ) = m a t h r m r n ( X × Y + Z ) mathrm{FMA}(X, Y, Z) = mathrm{rn}(X × Y + Z)mathrmFMA(X,Y,Z)=mathrmrn(X×Y+Z)少一次舍入的价值有多大？看一个只在浮点上成立、在实数上不成立的例子（binary32）：取A = 2 0 × 1.00000000000000000000001 A = 2^{0} × 1.00000000000000000000001A=20×1.00000000000000000000001，B = − 2 0 × 1.00000000000000000000010 B = -2^{0} × 1.00000000000000000000010B=−2
