FP8 的「不可能三角」被打破了DeepGEMM 深度解析里的精度与速度博弈【免费下载链接】DeepGEMMDeepGEMM: clean and efficient BLAS kernel library on GPU项目地址: https://gitcode.com/GitHub_Trending/de/DeepGEMM2025 年 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 或 128SM100粒度越细对激活值动态范围的适配越好但缩放因子的存储和搬运开销也越大。这是 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.intint32见 deep_gemm/utils/math.py 中的pack_ue8m0_to_int与ceil_to_ue8m0。于是每个 K 位置块的缩放信息只占 1 字节MN 方向连续排列既省带宽又满足 TMA 的 16 字节对齐要求。这套「幂缩放 打包」设计的巧妙之处在于它把精度问题的大部分负担从「乘法的舍入」转移到了「量化器的选择」上。用户侧量化时round_sfTrue四舍五入到幂内核侧消费时缩放是精确的误差被压缩到「量化」这一个环节边界清晰、可审计。补偿机制二硬件级协同让 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_castuint64_t(static_castuint32_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) 1MN 连续、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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
FP8 的「不可能三角」被打破了DeepGEMM 深度解析里的精度与速度博弈【免费下载链接】DeepGEMMDeepGEMM: clean and efficient BLAS kernel library on GPU项目地址: https://gitcode.com/GitHub_Trending/de/DeepGEMM2025 年 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 或 128SM100粒度越细对激活值动态范围的适配越好但缩放因子的存储和搬运开销也越大。这是 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.intint32见 deep_gemm/utils/math.py 中的pack_ue8m0_to_int与ceil_to_ue8m0。于是每个 K 位置块的缩放信息只占 1 字节MN 方向连续排列既省带宽又满足 TMA 的 16 字节对齐要求。这套「幂缩放 打包」设计的巧妙之处在于它把精度问题的大部分负担从「乘法的舍入」转移到了「量化器的选择」上。用户侧量化时round_sfTrue四舍五入到幂内核侧消费时缩放是精确的误差被压缩到「量化」这一个环节边界清晰、可审计。补偿机制二硬件级协同让 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_castuint64_t(static_castuint32_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) 1MN 连续、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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
