企业官网建设流程全解析

FP8 的「不可能三角」被打破了？DeepGEMM 深度解析里的精度与速度博弈

【免费下载链接】DeepGEMMDeepGEMM: clean and efficient BLAS kernel library on GPU项目地址: https://gitcode.com/GitHub_Trending/de/DeepGEMM

2025 年 2 月，DeepSeek 开源周第三天放出的 DeepGEMM，用一份仅数百行的 CUDA 代码库把 FP8 矩阵乘法的讨论从「能不能用」重新拉回「怎么才能用得更好」。社区情报里频繁出现三个数字——1350+ TFLOPS、1550 TFLOPS、比 cuBLASLt 快 2.7 倍——它们指向同一个事实：FP8 的算力红利是真实存在的，而它的代价（精度损失）也从未消失。本文不打算复述「300 行代码吊打英伟达」的叙事，而是回到仓库源码，逐层拆解 DeepGEMM 为 FP8 付出的「精度补偿成本」，以及这些成本最终如何被硬件与软件协同消化。读完你会得到一个更冷静的结论：所谓「不可能三角」并没有被物理性地打破，它只是被重新定价了。

FP8 的精度陷阱：为什么低精度总被质疑

FP8 的质疑史几乎与它本身一样长。E4M3 格式只有 1 位符号、4 位指数、3 位尾数，动态范围约 ±448，尾数精度与 FP16 相差一个数量级。做一次乘法还好，累计 K 次点积求和，截断误差会被放大 K 倍——这正是大模型训练早期宁可扛着显存压力用 FP16/BF16，也不肯碰 FP8 的原因。

但问题不止于「位数少」。工业界的经典疑虑有三个层面：

• 范围失配：激活值和权重值的量级差异很大，一个全局缩放因子无法同时覆盖两者；
• 累加漂移：即使输入被量化到位，FP32 之外的累加仍会把误差滚雪球；
• 格式不统一：各家库的量化粒度、缩放因子格式千差万别，换个库精度就变，工程上无法接受。

DeepGEMM 的选择是承认「误差无法消除」，然后把它拆成一个个可量化、可对冲、可验证的小问题。在 docs/scaling-factor-format.md 中，量化方案被显式定义为一个三元组recipe = (gran_m, gran_n, gran_k)，即缩放因子（SF）的存储粒度：A 矩阵每gran_m × gran_k一块共享一个缩放值，B 矩阵每gran_n × gran_k一块共享一个。gran_k 支持 32 或 128（SM100），粒度越细，对激活值动态范围的适配越好，但缩放因子的存储和搬运开销也越大。这是 DeepGEMM 对「精度 vs 吞吐」做的第一笔显式交易。

补偿机制一：把缩放因子做成「零误差」的幂

细粒度量化解决了范围失配，但缩放因子本身也会引入误差——普通的 FP32 缩放因子与数据相乘后再量化，会产生二次舍入。DeepGEMM 在这里选了最「省事」也最「极端」的路线：缩放因子必须是精确的 2 的幂。

在 docs/scaling-factor-format.md 的第 1.2 节，这条约束被写成硬性校验：

每个 float32 SF 值必须为精确的 2 的幂：bit pattern[0][8-bit exponent][23 mantissa bits = 0]，符号位与尾数必须为 0，设备端断言(value & 0x807fffffu) == 0。

2 的幂缩放意味着什么？被量化数据与缩放因子的乘法退化为指数相加（硬件上等价于移位），在 FP32/BF16 域内都是精确操作，零舍入误差。再进一步，SM100 上这些幂次指数被打包成 NVIDIA 的 UE8M0 格式——8 位无符号指数，4 个装进一个torch.int（int32），见 deep_gemm/utils/math.py 中的pack_ue8m0_to_int与ceil_to_ue8m0。于是每个 K 位置块的缩放信息只占 1 字节，MN 方向连续排列，既省带宽又满足 TMA 的 16 字节对齐要求。

这套「幂缩放 + 打包」设计的巧妙之处在于：它把精度问题的大部分负担从「乘法的舍入」转移到了「量化器的选择」上。用户侧量化时round_sf=True（四舍五入到幂），内核侧消费时缩放是精确的，误差被压缩到「量化」这一个环节，边界清晰、可审计。

补偿机制二：硬件级协同，让 MMA 指令自己吃掉缩放

细粒度缩放要做到「无损提速」，关键是不让缩放因子的加载与应用变成软件开销。SM100 的答案是硬件原生的 block-scaled MMA。在 deep_gemm/include/deep_gemm/mma/sm100.cuh 中：

CUTLASS_DEVICE uint64_t make_runtime_instr_desc_with_sf_id( cute::UMMA::InstrDescriptorBlockScaled desc, const uint32_t& sfa_id, const uint32_t& sfb_id) { desc.a_sf_id_ = sfa_id, desc.b_sf_id_ = sfb_id; return static_cast<uint64_t>(static_cast<uint32_t>(desc)) << 32; }

缩放因子的索引（a_sf_id_/b_sf_id_）被直接写进 UMMA 指令描述符，Tensor Core 在执行tcgen05.mma时按块读取 SF 并完成动态缩放，无需软件逐元素处理。SF 本身则通过make_sf_desc以 UTCCP 布局进驻共享内存——Atom size: 8 x 128 bits，专门为「每 128 K 子块一个缩放」设计。SM90 没有这套硬件，于是 DeepGEMM 只支持 FP32 格式的缩放因子、走软件路径（Kernel1D2D），性能与精度特性也随之不同——见 csrc/apis/gemm.hpp 中按arch_major分发的分支。

这解释了社区情报里反复出现的「硬件级协同优化」到底指什么：不是玄学，是把缩放因子当作 MMA 指令的一等公民，让 8 位数据的缩放成本在硬件里归零。

补偿机制三：布局契约与 JIT，把「对齐」变成可维护的约束

软件库层面，DeepGEMM 的补偿方式是「把正确性做进类型系统」。SF 的布局不是自由格式，而是有严格 stride 契约的 MN-major TMA 对齐布局，由 csrc/utils/layout.hpp 的check_sf_layout在主机侧逐项断言：stride(-2) == 1（MN 连续）、stride(-1) == get_tma_aligned_size(mn, element_size)（16 字节对齐）。凡是布局不满足契约的输入，调用时直接DG_HOST_ASSERT失败，而不是让内核在显存里读错数据、静默产出错误结果。

对应地，docs/scaling-factor-format.md 的 Section 3 定义了预变换 SF 的完整契约：[·, mn, ceil_div(k, gran_k*4)]的 int32 张量、MN-major、TMA 对齐。权重侧可以「变换一次、缓存复用」，激活侧则直接从 cast 内核产出打包好的 SF，后续 GEMM 调用只做校验、不再启动变换内核——消除重复开销的同时，也把「布局错误」挡在运行时之前。配合 csrc/runtime/jit.hpp 的 DeepJIT 运行时编译（安装零 CUDA 编译、按 shape 即时生成内核），这一切约束与调优都发生在同一个轻量代码库里，这也正是「简洁」与「正确」能共存的原因。

精度与吞吐的平衡点：源码里的数字说了算

质疑「FP8 够不够准」最好的回应不是口号，而是测试代码里的阈值。在 tests/test_fp8_fp4.py 中，QuantConfig.max_diff()定义了不同量化组合与 BF16 参考结果的允许偏差：

def max_diff(self) -> float: if self.is_fp4_a and self.is_fp4_b: return 0.02 if self.is_fp4_a or self.is_fp4_b: return 0.01 return 0.001

纯 FP8 组合（E4M3 双端）要求与 BF16 参考的差异小于 0.001，混合 FP4 放宽到 0.01~0.02——这是对「误差可控」的量化承诺。更值得注意的是同一文件里的两组测试：

• 位级确定性：同一输入跑 20 次，torch.equal断言输出逐位一致——精度之外，工程上更怕的是「同样代码两次跑出不同结果」，这直接服务训练复现与 CUDA Graph 推理；
• FP4/FP8 等价性：calc_diff(equivalent_d, equivalent_fp8_d) < 1e-14，即 FP4 结果与「把 FP4 数据转成 FP8 再算」的结果几乎逐位一致，说明误差的主源在量化而非内核计算路径。

而 tests/test_bf16.py 对 BF16 路径的阈值是< 1e-5，与 cuBLASLt 的平均加速比也会打印出来供横向对比。把两组阈值放在一起看，工业场景的「平衡点选择」其实有了清晰的操作指南：

• 对精度极其敏感、无法容忍量化偏差的层：走bf16_gemm_*，误差 < 1e-5，代价是吞吐折半量级；
• 主流前向与梯度计算（权重/激活均为 E4M3）：走fp8_gemm_*，误差 < 1e-3，换取接近硬件峰值的吞吐——README 中记录的 H800 上 1550 TFLOPS 正属于这条路径；
• 最激进的显存/带宽优化（FP4 专家权重，如 Mega MoE）：接受 1e-2 量级误差，换来的是一半的权重体积和两倍于 FP8 的单核计算密度。

DeepSeek-V3 的实践（约 280 万 H800 GPU 小时完成训练）表明，在 MoE 架构 + FP32 累加 + 细粒度量化 + 硬件级动态缩放的组合下，FP8 的误差被压在收敛可接受范围内，同时把训练成本推到此前不可想象的量级。这也是 DeepGEMM 与 DeepEP、FlashMLA 等一起被开源的核心原因——误差不是被消灭了，而是被分摊到了「可控、可验证、可复现」的工程环节里。

结论：被打破的是「盲目」，不是「物理」

回到标题的问题：FP8 的「不可能三角」被打破了吗？从仓库源码能得出的准确答案是——DeepGEMM 没有让 FP8 同时做到「无限精度 + 极限吞吐 + 完全通用」，它做的是三件更务实的事：把缩放因子变成零误差的幂并交给硬件（deep_gemm/include/deep_gemm/mma/sm100.cuh）、把布局契约变成运行时的硬性校验（csrc/utils/layout.hpp）、把误差边界变成测试里的明确数字（tests/test_fp8_fp4.py）。「不可能三角」没有消失，只是从「玄学判断」变成了「工程定价」：每一档精度阈值都有对应的吞吐与成本，选择权被明明白白地交还给了使用者。

对于部署大模型的工程师，真正的启示是：别再问「FP8 能不能用」，要问「我的量化粒度、缩放格式、累加精度、验证阈值各是多少」。DeepGEMM 用一份整洁的源码把这些问题全部显式化了——这或许比任何性能数字都更有价值。

【免费下载链接】DeepGEMMDeepGEMM: clean and efficient BLAS kernel library on GPU项目地址: https://gitcode.com/GitHub_Trending/de/DeepGEMM