Deploying DeepSeek on GB200 NVL72 with PD and Large Scale EP (Part II): 3.8x Prefill, 4.8x Decode Throughput

发表时间: 2025-09 · Blog post by The SGLang Team (lmsys.org)

原文: https://lmsys.org/blog/2025-09-25-gb200-part-2

作者/机构：The SGLang Team

速读

一句话结论 SGLang 团队通过引入 FP8 注意力机制、NVFP4 混合专家计算、细粒度通信计算重叠以及基于卸载的缩容策略，在 GB200 NVL72 硬件上部署 DeepSeek V3/R1 模型，实现了每秒每 GPU 26156 个输入 token 和 13386 个输出 token 的吞吐量，相比 H100 环境分别提升了 3.8 倍和 4.8 倍。

要解决什么问题 在大规模混合专家模型（如 DeepSeek V3/R1）的推理部署中，原有的常规精度和并行策略在顶级硬件上面临多重底层物理机制的卡点。首先是显存访问压力与带宽瓶颈，在解码阶段，传统的 BF16 精度键值缓存（KV cache）会占用大量显存带宽，导致注意力计算严重受限于访存速度而非算力；同时，通用矩阵乘法（GEMM）在使用 FP8 精度时，依然面临极高的显存带宽压力，限制了计算单元的利用率。其次是通信开销与显存容量的矛盾，在进行大规模专家并行（EP）时，跨节点或跨设备的 token 分发会产生庞大的网络通信流量，且模型权重本身占据了大量显存，导致留给 KV cache 的空间被严重挤压，进而限制了系统能够支持的并发批处理大小和上下文长度。最后是计算与通信的重叠效率问题，随着新一代硬件（如 GB200）通信带宽的显著提升，原有针对旧硬件设计的粗粒度双批次重叠策略已无法完全掩盖通信延迟，系统经常需要浪费时间等待最慢的计算节点，导致整体并行度无法进一步提升。

怎么做的 为了彻底绕开上述访存与通信卡点，该方法在精度压缩、显存调度和底层算子三个维度进行了深度重构。第一，引入 FP8 注意力机制，将 KV cache 的数据格式从 BF16 降至 FP8，这直接将解码阶段的显存访问压力减半，并允许调用更快的张量核心指令；更重要的是，单 token 占用的显存减少使得系统能容纳更多 KV cache，从而支持更大的批处理大小以提升整体效率。第二，全面采用 NVFP4 GEMM 替换经典的 FP8 GEMM，不仅进一步降低了矩阵乘法的显存带宽压力并激活了更强的 FP4 张量核心，还在 DeepEP 库中将 NVFP4 量化与 token 分发过程融合，使网络通信流量直接减半；此外，除了混合专家层，注意力机制的输出投影层也采用了 NVFP4 量化，但为了保底性能，其内部的 q_b_proj 投影仍保持 FP8 精度。第三，设计了基于卸载的缩容策略，当 GPU 显存不足以支撑庞大的专家并行规模时，利用 GB200 节点内 CPU 与 GPU 之间高达 900GB/s 的双向高带宽，配合预取机制将部分权重卸载到主机内存中；这种做法减少了单个预取实例所需的 GPU 数量，缩小了显存爆炸半径，并通过减少通信开销对冲了部分计算降速，同时缓解了等待最慢节点的长尾效应。第四，重构了细粒度的计算通信重叠机制，将合并通信与下投影矩阵乘法以及共享专家计算进行重叠；在底层实现上，利用张量内存访问（TMA）存储提交后足够多步的等待机制，配合带有释放语义的原子指令，以及 PTX 级别的 cp.async.bulk.wait_group 异步等待指令，实现了极致的流水线掩盖。在算子工程层面，系统集成了 NVIDIA Blackwell DeepGEMM 用于预填阶段的注意力计算，并基于 CuTe DSL 和 CUTLASS 框架分别定制了针对 NVFP4 的解码与预填 GEMM 算子，通过持久化分块调度和 warp 级特化技术榨干硬件性能。

效果如何 实验在 GB200 NVL72 硬件上进行，测试模型为 DeepSeek V3/R1，输入序列长度统一设定为 2000 个 token。解码阶段采用 48 个计算节点的大规模专家并行，预填阶段高精度模式使用 4 个节点，低精度模式使用 2 个节点；批处理大小根据 KV cache 容量动态拉满（例如 4K 序列长度下设为 768，2K 序列长度下设为 1408）。对比基线为该团队此前在 H100 硬件上采用传统精度（BF16 注意力配合 FP8 混合专家）的大规模专家并行路线。量化结果显示，在极致的低精度配置（FP8 注意力配合 NVFP4 混合专家及输出投影）下，预填和解码吞吐量分别达到每秒每 GPU 26156 个输入 token 和 13386 个输出 token，相比 H100 基线分别实现了 3.8 倍和 4.8 倍的绝对加速。在微观算子层面，FP8 注意力算子和 NVFP4 GEMM 算子相比原精度版本分别获得了最高 1.8 倍和 1.9 倍的提速。在代价与局限性方面，作者承认训练后量化不可避免地会带来信息损失，但由于 NVFP4 采用了极小的块大小（16）并保留 FP8 作为缩放因子，实验证明其精度下降微乎其微，与官方检查点保持一致。此外，作者也指出当前端到端的性能差异并不完全归功于精度降低，不同代码路径的辅助算子优化程度不一，且降低批处理大小会带来明显的性能惩罚（例如将 2K 序列的批处理大小从 1408 降至 768 会导致约十分之一的性能损失），未来仍需引入多 token 预测的重叠调度器来弥补小批次下的解码性能。

主要贡献

本文致力于在目前最强大的深度学习硬件之一 GB200 NVL72 上，优化 DeepSeek V3/R1 模型的推理性能。核心研究目标是通过引入低精度计算与系统级优化，解决大规模模型推理中的显存与通信瓶颈。

其主要创新点与贡献如下：

  1. 极致的吞吐量提升：在采用 FP8 注意力机制（Attention）、NVFP4 混合专家（MoE）、大规模专家并行（EP）以及预填充-解码分离（Prefill-Decode Disaggregation）等优化后，对于 2000-token 的输入序列，SGLang 在每个 GPU 上实现了预填充 26,156 tokens/s 和解码 13,386 tokens/s 的吞吐量。与 H100 的设置相比，分别实现了 3.8 倍和 4.8 倍的加速。
  2. 传统精度下的高性能保障：即使在使用传统精度（注意力机制使用 BF16，GEMM 使用 FP8）的情况下，SGLang 依然能够实现预填充 18,471 tokens/s 和解码 9,087 tokens/s 的高吞吐量。
  3. 低精度的显著收益与无损精度：与原始精度版本相比，使用 FP8 注意力机制和 NVFP4 GEMM 内核分别带来了最高 1.8 倍和 1.9 倍的性能提升，且这种极低精度的量化（FP8/NVFP4）带来的模型准确率下降微乎其微（negligible accuracy degradation）。

方法细节

为了在 GB200 NVL72 上最大化 DeepSeek 的推理性能，系统在宏观策略和底层内核级别实施了详尽的优化。

FP8 注意力机制（FP8 Attention）
在传统的 BF16 精度之外，系统现在支持在注意力机制中使用 FP8 精度来存储 KV 缓存。因为采用了 FP8 数据格式，这大致将解码（decode）阶段的内存访问压力减半，并且允许硬件使用更快的 Tensor Core 指令，从而直接加速了解码注意力内核的执行。此外，该设计还使得固定大小的显存中能够容纳更多的 KV 缓存 tokens，进而支持更长的序列和更大的批处理大小（batch size），而更大的批处理大小进一步提升了系统的整体执行效率。

NVFP4 通用矩阵乘法（NVFP4 GEMM）
与经典的 FP8 GEMM 相比，采用全新的 NVFP4 精度不仅降低了 GEMM 操作的内存带宽压力，还允许系统利用 GB200 上更强大的 FP4 Tensor Core 进行计算。其次，通过将所需的网络通信流量减半，该方法大幅加速了 token 的分发过程。最后，它显著降低了权重的内存消耗，这使得系统能够缩小部署规模（scaling down）或者为 KV 缓存腾出更多空间。除了 MoE（混合专家）的专家层在 NVFP4 精度下执行外，注意力机制中的输出投影（output projection）GEMM 也可以选择性地量化为 NVFP4。与 NVIDIA 官方的检查点不同，为了进一步提升性能，系统将 q_b_proj 在 FP8 而不是 BF16 精度下执行。

通过卸载实现规模缩减（Scaling Down by Offloading）
除了向上扩展（scaling up），系统还支持缩小专家并行（EP）的规模。当设备显存不足时，系统利用 GB200 在 CPU 和 GPU 之间的高速双向带宽（900GB/s），通过预取（prefetching）技术将权重卸载（offload）到主机内存中。因为减少了跨节点的通信开销，当这种通信收益超过计算变慢的代价时，就能带来整体性能的提升。因此，最佳的缩放比例不仅与模型配置有关，还与当前使用的计算和通信内核密切相关。这种方法也最小化了故障的爆炸半径（explosion radius），因为每个预填充实例使用的 GPU 数量更少。最后，这种设计还可以减少等待集群中最慢 rank 所浪费的时间。

计算与通信的细粒度重叠（Computation Communication Overlap）
鉴于 GB200 通信带宽的显著增加，以往硬件中使用的双批次（two-batch）重叠策略可能不再是最合适的方案，因此系统采用了一种细粒度的重叠方法。为了简化实现，系统将合并通信（combine communication）与向下 GEMM（down GEMM）以及共享专家（shared experts）进行了重叠。在 GEMM 中实现信号传递（signaling）时，系统在 TMA 存储等待（TMA store wait）之后使用了带有释放语义（release semantics）的原子指令，该等待操作发生在 TMA 存储提交（TMA store commit）之后的足够多步之后。此外，系统使用了 cp.async.bulk.wait_group PTX 指令，该指令系列通常用于 tma_store_wait 或等效过程中，但此处去除了 .read 后缀以适配特定需求。

内核级优化与集成
在底层内核级别，系统在 SGLang 中集成或深度优化了以下具体内核：

  • 用于预填充注意力的 NVIDIA Blackwell DeepGEMM：NVIDIA Blackwell DeepGEMM 是一个统一的内核，在预填充和解码中均实现了高性能。因此，除了在解码中使用它之外，系统还将该内核无缝集成到了预填充的代码路径中。
  • 用于 NVFP4 解码的 FlashInfer Blackwell CuTe DSL GEMM：该内核利用 CuTe DSL 实现了在 NVFP4 精度下具有掩码布局（masked layout）的 GEMM。它利用了张量内存访问（TMA）和 tcgen05.mma 指令（包括 2CTA MMA）来实现高效计算，同时还使用了持久化分块调度（persistent tile scheduling）和 warp 特化（warp specialization）技术。
  • 用于 NVFP4 预填充的 FlashInfer Blackwell CUTLASS GEMM：该模块支持多种数据类型，并通过 CUTLASS 实现。其应用的优化策略与 CuTe DSL 版本类似。该内核专为高吞吐量工作负载设计，特别适合预填充阶段。
  • 用于 BF16 KV 缓存预填充的 Flash Attention CuTe：与上述 GEMM 类似，该内核在 CuTe DSL 框架中编写，在预填充期间为多头注意力（MHA）实现了极高的性能。
  • 用于解码和 FP8 KV 缓存预填充的 FlashInfer Blackwell TensorRT-LLM Attention：该内核利用了基于集群启动控制（cluster launch control）的持久化调度器，有效地隐藏了内核的序言（prologue）和结语（epilogue）。它还实现了计算与内存加载之间更好的重叠。该内核同时支持 BF16 和 FP8 精度。
  • 在 DeepEP 中融合 NVFP4：DeepEP 在分发 token 之前可以选择性地对它们进行量化，因此 NVFP4 量化与原始的 FP8 操作一起被融合到 DeepEP 中，从而将所需的网络流量减半。
  • 较小内核的优化与融合（Smaller Kernels）：首先，对其他小内核（如量化和拼接内核）进行了深度优化和融合。其次，系统还优化了 FlashInfer 中的 MLA RoPE 量化内核。最后，系统还稍微优化了 TensorRT-LLM 中位于 FlashInfer 的几个内核作为原型，这带来了 5% 的端到端加速以及最高 2.5 倍的单内核加速。

实验环境

  • 硬件配置：NVIDIA GB200 NVL72 集群（对比基线为 H100 集群）。充分利用了 GB200 架构中 CPU 与 GPU 之间 900GB/s 的双向带宽。
  • 软件配置：SGLang 推理系统。实验中使用了早期访问版本（early access version）的 CuTe DSL，因为所需的 bug 修复尚未公开发布。
  • 模型配置：DeepSeek V3 / R1 模型。
  • 部署与参数设置：采用预填充-解码分离架构。对于解码阶段，使用 48 个 ranks（即大规模专家并行 EP）；对于预填充阶段，高精度模式下每个实例使用 4 个 ranks，低精度模式下使用 2 个 ranks。输入序列长度设定为 2000 tokens。数据保持分布内（in-distribution）。批处理大小（Batch size）的选择以大致填满 KV 缓存为准（例如 4k ISL 时为 768，2k ISL 时为 1408），若降低批处理大小（如 2k ISL 时从 1408 降至 768），性能会下降约十分之一。

实验结果

端到端性能评估

  • 实验内容：在 GB200 NVL72 上评估 DeepSeek 在 SGLang 中的端到端吞吐量，对比原始精度（BF16 注意力 + FP8 MoE）和降低精度（FP8 注意力 + NVFP4 MoE 及输出投影 GEMM）的差异，并与 H100 基线进行比较。
  • 实验结果：在 2000-token 的输入序列下，降低精度配置实现了预填充 26,156 tokens/s/GPU，解码 13,386 tokens/s/GPU，相比 H100 分别加速了 3.8 倍和 4.8 倍。在原始精度配置下，依然达到预填充 18,471 tokens/s，解码 9,087 tokens/s。
  • 分析结论：这种显著的加速主要归因于：1）降低精度减少了计算和内存访问开销，并允许更大的批处理大小；2）更快的内核（如集成的 DeepGEMM 和 NVFP4 专用内核）占据了端到端时间的重要部分；3）多种优化手段（如重叠、卸载、小内核融合等）产生了乘数效应。此外，高低精度代码路径的性能差异不仅源于精度本身，还因为它们采用了不同的辅助内核和调度策略。
  • 图表引用：
    端到端性能对比图

低精度内核加速效果分析

  • 实验内容：单独考察注意力内核和 GEMM 内核（包括 MoE 中的 gate-up GEMM、down GEMM 以及注意力中的输出投影 GEMM）在从标准精度切换到低精度后的速度变化。
  • 实验结果：降低精度极大提升了相关内核的执行速度。在测试用例中，注意力内核加速了 1.8 倍，GEMM 内核加速最高达到了 1.9 倍。
  • 分析结论：除了直观的内核级执行时间缩短，降低精度带来的另一个无法直接从内核视角看到的提升是 KV 缓存 token 数量的增加，这直接导致了批处理大小的扩大，从而进一步提升了整体吞吐性能。
  • 图表引用：
    内核级加速效果图

模型精度损失评估

  • 实验内容：评估采用 NVFP4 进行训练后量化（PTQ）是否会导致模型性能和准确率的下降。
  • 分析与结果：理论上，NVFP4 选择较小的块大小（block size 为 16）并使用 FP8 作为缩放因子（scaling factor）数据类型，使其能够以尽可能少的损失表示原始信息。实验观察表明，评估结果与 NVIDIA 针对 NVFP4 的官方检查点一致，量化带来的模型准确率变化极小（tiny accuracy changes），几乎可以忽略不计。
  • 图表引用：
    模型准确率评估图

结论

本文通过结合 FP8/NVFP4 极致低精度计算、大规模专家并行、预填充-解码分离以及一系列底层内核与通信调度优化，在 GB200 NVL72 硬件上实现了 DeepSeek V3/R1 模型推理性能的巨大飞跃，同时保持了模型的准确率。

对于未来的工作，研究团队提出了以下展望：

  1. 带重叠调度器的多 Token 预测（MTP）：这将非常有利于解码阶段的性能提升，特别是在内核批处理较小或注意力内存访问压力较大时。
  2. 进一步的内核优化：部分内核仍有改进空间，以求完全榨干硬件潜力。
  3. 支持更多模型：计划将这些优化策略应用于其他强大的模型，如 Kimi-K2、Qwen 和 GLM 4.5。
  4. OME 集成与简化部署：利用 OME（用于企业级 LLM 管理和服务的 Kubernetes operator）来简化模型的部署和日常运维操作。