vLLM 支持 NVIDIA Vera Rubin NVL72，吞吐量达 GB200 NVL72 的 7.8 倍

vLLM 现已支持 Vera Rubin NVL72！

NVIDIA Vera Rubin 是专为智能体推理打造的下一代平台。自 Vera Rubin NVL72 发布以来，Inferact、NVIDIA、Red Hat 和 vLLM 社区一直在将 vLLM 移植到 Vera Rubin NVL72 上，如今 vLLM 已可在 Vera Rubin NVL72 上运行，具备每日容器构建，并支持来自 DeepSeek、Moonshot AI、Z.ai 和 MiniMax 的模型。

本文是对当前进展的初步介绍，以下是迄今为止这项工作的一些亮点：

• Vera Rubin NVL72 硬件： NVFP4 FLOPS 达到 GB200 NVL72 的 5 倍，HBM 带宽约为其 2.4 倍，双向 NVLink 带宽为其 1.7 倍，softmax 的指数运算速度提升 2-4 倍。
• Day-0 支持： Rubin 基于 Blackwell 的架构家族构建，因此 vLLM 的 Blackwell 内核与 Rubin 兼容。得益于此，vLLM 已支持在 Rubin 上运行 DeepSeek、Kimi、GLM 和 MiniMax 等多种模型。
• Rubin 调优内核： 通过 FlashInfer 0.7.0，vLLM 获得了针对 Rubin 调优的 attention、GEMM 和 MoE 内核。我们还针对 Rubin 调优了 MiniMax 稀疏注意力（MSA）prefill 内核。
• 局部性感知 MoE： 为充分利用 Rubin 更高的 HBM 带宽，我们利用 CUDA 13.4 的局部性域（locality domains）来拆分 MoE 权重。这使得 SM 只能从距离自己最近的内存中读取权重。
• 早期性能： 早期结果已显示出 vLLM 令人瞩目的提升：在 AgentX 上、交互性匹配的情况下，每 GPU 吞吐量达到 GB200 NVL72 的 7.8 倍；在 MLPerf 中，VLM 吞吐量相比 GB300 NVL72 最高提升 3.7 倍。这仅仅是开始；随着优化持续进行，我们预计还会看到更多性能提升。

Rubin 为推理带来的变化

图 1. NVIDIA Vera Rubin NVL72 与 GB200 NVL72 的每 GPU 对比。将鼠标悬停在某项指标上可高亮其在 GPU 中的对应部分；“Show table”会列出所有数值。来源：NVIDIA Vera Rubin NVL72 和 GB200 NVL72 规格页面，以及 NVIDIA Rubin 开发者博客。

Vera Rubin 平台

Vera Rubin 平台通过对机架组件进行极致的协同设计来实现出色性能——它针对智能体 AI 工作负载提供了五个全新、独特、专为机架规模打造的系统：Vera Rubin NVL72、Vera CPU 机架、Groq 3 LPX、Spectrum-6 SPX 和 BlueField-4 STX Storage。

单台 Vera Rubin NVL72 提供的 NVFP4 推理 FLOPS 是 GB200 NVL72 的 5 倍，内存带宽为其 2.4 倍。在 scale-up 网络方面，第六代 NVLink 提供的带宽最高可达 Blackwell 的 1.7 倍，在生产级智能体服务场景中带来显著更好的用户体验。

Softmax。 值得注意的是，Rubin 还提升了 softmax 性能，这是 LLM 注意力中的核心操作。Rubin 提高了指数运算吞吐量，其中 FP32 吞吐量为 NVIDIA GB200 的 2 倍，BF16/FP16 吞吐量为 4 倍，帮助 softmax 跟上更快的矩阵运算。

内存。 Rubin 中的 HBM 已从 HBM3e 升级到 HBM4，相比 GB200 NVL72 提供最高 2.4 倍的带宽。结合其更强大的计算能力，Rubin GPU 加速了 GEMM、MoE（混合专家）和注意力等关键 LLM 推理操作，并带来更高的整体吞吐量和更低的解码延迟（详见下文）。

网络。 Rubin 平台上的 GPU 间网络也得到了大幅改进。第六代 NVLink 提供的网络带宽比上一代高 1.7 倍。这将加速集合通信（例如 AllReduce 和 All2all）及其他通信操作，提升大规模 LLM 推理的速度和可扩展性（例如 prefill/decode 分离、宽专家并行）。

vLLM 对 Rubin 的支持状态

vLLM 社区在 Rubin 公开发布后便立即着手进行 Rubin 适配工作。作为 vLLM 社区的重要成员，来自 NVIDIA、Inferact 和 Red Hat 的工程师们通力合作并做出贡献，确保所有用户都能轻松地在 Rubin 平台上部署任意模型。

在本节中，我们重点介绍 vLLM 正在进行的 Rubin 专项支持，即利用局部性域（locality domain），以及我们为在 Rubin 硬件上开箱即用所做的易用性改进。

局部性域支持

自 Ampere 以来，NVIDIA GPU 就具备非均匀全局内存访问特性。NVIDIA CUDA 13.4 中的局部性域功能允许应用程序通过将计算和数据放在同一局部性域内，充分利用非均匀全局内存访问。SM 访问自身局部性域内的全局内存时，相比访问其他域中的 HBM，可获得更高带宽和更低延迟。借助Green Contexts 和 CUDA 流，我们可以在每个局部性域中各启动一个内核，使每个内核都能访问本地内存。该功能主要加速受内存带宽限制的工作负载，例如 MoE decode。局部性域仍处于积极的设计和开发之中。在本节中，我们以 MoE decode 为例进行深入探讨。

图 3. 在两个局部性域上进行 MoE 前向计算的 Split-N。权重 W 按列在 N/2 处拆分，每个域的 SM 只读取自身 HBM 中的那一半 W，因此每个域使用其本地内存带宽。输入 X 和输出 C 跨越两个域。使用暂停、上一个/下一个或步骤标签来逐步查看。

MoE decode 受限于从 HBM 读取权重，因此我们的目标是优化跨局部性域的内存吞吐量。作为对 Rubin 新局部性域功能的初步探索，我们在 FC1 和 FC2 中都采用 split-N 策略，如图 3 所示。我们按列对权重进行分片，将每个分片放入各自局部性域的全局内存中，并将每个域的 SM 限制在其本地分片上。这消除了大部分跨域内存访问，从而提升内核性能并节省功耗。由于 decode 期间激活内存相对较小，将其跨两个内存域保持非本地化所带来的开销极小。

SM 并不总能被划分为均等的域。在尝试创建均等分区时，局部性域创建默认不会包含那些 SM。为使两个分区拥有相等的 SM 数量，我们需要在创建域时启用 cudaDevSmResourceGroupBackfill（回填模式）（更多细节请参阅官方局部性域文档）。在我们的性能研究中，我们同时纳入了默认模式（两个域总共仅使用 200 个 SM）和回填模式（全部 212 个 SM 均被使用）。

图 4 比较了在不同并行策略下，开启和关闭局部性域时 MoE 层前向时间（FC1 + FC2）的初步结果。我们以 MiniMax M3 MoE 的 shape 为例。启用局部性域后，在小 token 前向场景下，我们能够稳定获得平均 1.2 倍的加速。对于其他 TP 和 EP 服务策略，这一趋势大致相同。即使在默认模式下（仅使用 212 个 SM 中的 200 个），启用局部性域也能带来类似的收益。主要原因是，在小 token 解码中，前向传播主要受权重加载主导，而局部性域能够实现更高的 HBM 吞吐量。这些早期结果只是一个起点，在 Rubin 上仍有进一步调优和优化的空间，以最大化本地化带来的性能收益。

图 4. Rubin 上 MiniMax M3 MoE 层每 rank 的 FC1 + FC2 初步延迟，非本地化 vs 本地化（越低越好），每对上方标注加速比（非本地化延迟 ÷ 本地化延迟）。标签页切换并行策略（TP2、TP4、EP2、EP4）；开关在回填模式（全部 212 个 SM）和默认模式（212 个 SM 中的 200 个）之间切换。均衡路由；不包含通信时间。

Day-0 可用性

可用性始终是 vLLM 的首要优先级。截至今天，用户可以从 vLLM 的 Docker Hub 拉取并使用基于 CUDA 13.4 和 PyTorch 2.15 构建的 nightly 镜像，即 vllm/vllm-openai:cu134-nightly，用于 Rubin 硬件。

Blackwell 软件栈兼容性。Rubin 构建于 Blackwell 架构家族之上，并扩展了 tcgen05 tensor core 指令。它是一个新的 GPU 编译目标（sm107），但为 Blackwell 家族目标（sm100f）构建的 kernel 也可以在其上运行。在实践中，vLLM 的 Blackwell kernel，尤其是 GEMM 密集型的 kernel，如 attention 和 MoE，已经可以在 Rubin 上无需任何修改地运行。

每日容器构建。Rubin 的每日容器构建已经可用（#55953），这得益于 #53443 和 #54640 为 CUDA 13.4 上的 Rubin 构建路径提供支持，以及 #56545 和 #59288 对 Rubin 依赖项的更新。

模型覆盖。有了这些组件，vLLM 现在可以在 Rubin 上服务包括 DeepSeek、Kimi、GLM 和 MiniMax 在内的多种模型。

Rubin 调优 kernel

利用 Rubin 特定硬件能力和特性的 kernel 正逐步发布并上游到 kernel 库，如 FlashInfer、vLLM 的 MSA 分支（vllm-project/MSA）、Humming（vllm-project/humming）等。截至今天，vLLM 已集成了一些重要的 kernel 以实现最大化的 Rubin 性能，包括 dense NVFP4 或 MXFP4 GEMM、NVFP4 MoE、FP8 attention、FP8 MSA prefill 等。

性能

我们使用两个代表性基准评估了 vLLM 在 Vera Rubin NVL72 GPU 上的性能：SemiAnalysis AgentX（详见我们之前的文章）和 MLPerf Inference v6.1。

在 AgentX 上，vLLM 在 Vera Rubin NVL72 上运行 MiniMax M3，在匹配交互性下，吞吐量最高可达 NVIDIA GB200 的 7.84 倍，在 150 TPS 约束下吞吐量高出 5.18 倍。这是对该平台推理能力的非常早期的观察。随着我们获得更多 Vera Rubin NVL72 节点，我们将扩大测试范围并加速优化，随着这项工作的推进，预计将带来进一步的性能提升。

MLPerf Inference v6.1 轮次是首次将 vLLM 引入 NVIDIA Vera Rubin NVL72 平台的测试场。在视觉语言模型（VLM）基准测试中，通过 vLLM 作为后端推理引擎、Dynamo 作为前端路由器部署 Qwen3-VL-235B-A22B 模型，Vera Rubin NVL72 在离线、服务器和交互式场景下的吞吐量比 GB300 NVL72 高出最高 3.7 倍。有关已发布的 MLPerf Inference v6.1 结果的更多详情，请见 NVIDIA 的博客文章。

后续步骤

“罗马不是一天建成的”，打磨 Vera Rubin NVL72 GPU 上的可用性和性能将是一段充满兴奋的持续旅程。在不久的将来，作为社区共同努力，我们计划为 Rubin GPU 启用更多新功能，包括但不限于：

• 通过 FlashInfer 将 sm107 FlashInfer MegaMoE 集成到 vLLM 中。
• 全面启用 MoE 层的局部性域。
• 通过 PDL 和 Lamport Sync 发现并利用不同层或内核之间更多的重叠机会。
• 探索面向延迟敏感用例的巨型内核。
• 为 Kimi K3 优化 Rubin 上的 KDA 和 MLA 内核。
• 为 DeepSeek-V4.1-Flash 集成 Rubin CSA 和 HCA 内核。
• 完成并集成 Rubin MSA 解码内核。
• 通过 FlashInfer 将 CFT 计数写入 MoE all-to-all 内核集成到 vLLM 中。

致谢

这项工作是由 Inferact、NVIDIA、Red Hat 以及更广泛的 vLLM 社区共同协作完成的。我们特别感谢：

• NVIDIA，感谢其提供 Vera Rubin NVL72 的早期访问权限，并在整个开发过程中密切合作。
• Inferact 和 NVIDIA，感谢其领导 Rubin 协作、推动性能调优、集成 MSA 内核，并剖析局部性感知 MoE 的性能。
• NVIDIA 和 Red Hat，感谢其为 Rubin 设置并启用每日 Docker 构建。
• vLLM 社区，感谢其在整个过程中的持续支持和贡献。

附录：在 Vera Rubin NVL72 上运行 vLLM

展示 Rubin 的内核配置

vLLM 已经为 Rubin 集成了几个高度优化的内核。本节介绍相应的配置，以便在 Rubin GPU 上启用它们以获得最大性能。

• CuTe-DSL 密集 NVFP4 或 MXFP4 GEMM。对于 NVFP4/MXFP4 模型检查点默认开启，但你可以为 NVFP4 设置 --linear-backend flashinfer_cutedsl，或为 MXFP4 设置 --linear-backend flashinfer_cutlass，以确保其已启用。
• CuTe-DSL NVFP4 MoE。你可以通过 --moe-backend flashinfer_cutedsl 将其开启。
• 用于“batched”专家格式中 NVFP4 W4A4 MoE 的 CuTe-DSL 掩码分组 GEMM。这适用于 --enable-expert-parallel --data-parallel-size N --all2all-backend deepep_low_latency|nixl_ep 且 N>1 的部署。在这种情况下，--moe-backend auto|flashinfer_cutedsl 都会解析到此内核。
• 用于静态逐张量 FP8 W8A8 线性层的 CuTe-DSL FP8 BMM。默认开启，并且可以被 FlashInfer 自动调优器选中。
• Trtllm-gen FP8 注意力。要开启，你需要通过 --kv-cache-dtype fp8 使用 FP8 KV 缓存，或使用指定 FP8 KV 缓存的检查点，并同时设置 --attention-backend FLASHINFER|FLASHINFER_MLA。对于 DeepSeek 风格的 MLA 预填充，请同时添加 -ac.mla_prefill_backend=TRTLLM_RAGGED -ac.use_prefill_query_quantization=true。
• CuTe-DSL FP8 MSA 预填充。要开启，你需要通过 --kv-cache-dtype fp8 使用 FP8 KV 缓存，或使用指定 FP8 KV 缓存的检查点，并同时设置 --attention-config.minimax_m3_msa_decode_backend=cutlass。