DeepSeek-V4 on Day 0: From Fast Inference to Verified RL with SGLang and Miles
DeepSeek-V4 on Day 0: From Fast Inference to Verified RL with SGLang and Miles
发表时间: 2026-04 · Blog post by The SGLang and Miles Team (lmsys.org)
原文: https://lmsys.org/blog/2026-04-25-deepseek-v4
作者/机构:The SGLang and Miles Team (April 25, 2026)
速读
一句话结论
SGLang 和 Miles 团队在 DeepSeek-V4 发布首日,针对其混合稀疏注意力、流形约束超连接和 FP4 专家权重等全新架构,构建了首个支持全量推理与强化学习训练的开源系统栈,实现了接近百万上下文长度下的平稳高吞吐解码与稳定的强化学习训练。
要解决什么问题
DeepSeek-V4 的新架构引入了多个严重的系统级卡点。首先是前缀缓存失效:该模型每层混合使用滑动窗口注意力(SWA)和压缩注意力(C4 稀疏或 C128 稠密),导致系统必须在预填充、解码和投机解码阶段,强行维持三个异构 KV 池和两个压缩状态池的生命周期一致性,传统缓存假设彻底崩溃。其次是投机解码的调度瓶颈:混合注意力需要极其繁重的元数据(如页面索引、映射槽位、压缩计划),若在 CPU 调度流上急切准备这些数据,其发射开销会直接吞噬投机解码的收益。第三是显存与访存灾难:在 1M 上下文下,C4 稀疏注意力的索引器每步会产生 256K 个候选者,传统的全局排序耗时超过 100 微秒,直接阻塞流水线;而朴素的 KV 压缩操作需要 5 次 HBM 读写,导致压缩本身的开销甚至高于注意力计算;此外,C4 层绝大部分 KV 缓存处于非活跃状态,全量驻留 GPU 会造成极大的显存浪费。最后,在强化学习训练端,模型新增的随机算子(如 MoE 路由和索引器 Top-K)以及混合精度下的长轴 Softmax 累加,极易引发数值溢出和训练跑飞。
怎么做的
针对前缀缓存,作者提出了 ShadowRadix 机制。其核心思路是建立一个统一的“虚拟全 Token 槽位”基数树,将生命周期解耦:从虚拟槽向物理池(SWA、C4、C128)投射“影子”映射。系统为每个节点维护双计数器,当滑动窗口移出时,SWA 计数器归零并释放 SWA 物理槽,但节点及其 C4/C128 的压缩 KV 影子依然存活,供后续请求进行前缀匹配。嵌套的环形影子天然继承了这一特性,其物理地址定义为 $\text{swa\_page} \times \text{ring\_size} + \text{pos} \bmod \text{ring\_size}$,释放 SWA 页面会自动使内部的环失效。为了兼容投机解码,只需将环形缓冲区大小翻倍(如 C4 从 8 增至 16),即可确保被拒绝的草稿写入不会覆盖活跃窗口。针对调度瓶颈,采用图内元数据准备,将复杂的索引算术逻辑全部融合进 CUDA Graph,Python 端仅传入原始批次状态,彻底消除单步发射开销。针对显存与访存卡点,引入三个关键设计:一是 HiSparse 架构,利用 C4 层的稀疏性,将非活跃的压缩 KV 卸载到 CPU 内存池,GPU 仅保留小容量活跃工作集,通过 LRU 策略异步换入换出;二是 Flash Compressor,将 5 阶段的压缩流水线融合成单次片上处理,把 HBM 往返次数从 5 次砍到 2 次;三是 Lightning TopK,用 8 线程块集群的基数选择归约替代全局排序,各 CTA 构建局部 10 位基数直方图并归约出阈值,仅分散高于阈值的条目。在强化学习端,Megatron-LM 后端实现了全维度并行,其中上下文并行(CP)通过 All-gather 巧妙处理了 C4 压缩组的跨卡重叠偏移;同时引入路由重放(R3)和实验性的索引器重放机制来对齐随机算子,并将压缩器反向传播中容易因 BF16 累加产生大数吞小数偏差的 All-reduce 操作强制切换为 FP32,以保住数值稳定性。
效果如何
实验在 30K 截断的《红楼梦》长文本提示词下进行,硬件采用 B200(运行 1.6T Pro 模型,TP=8)和 H200(运行 285B Flash 模型,TP=4),对比基线为官方配方下的“其他开源引擎”(原文未具名)。在推理吞吐方面,得益于 ShadowRadix 和图内元数据优化,SGLang 在 4K 到 900K 上下文长度内吞吐量几乎保持平坦,衰减不到 10%(B200 上从 199 降至 180 tokens/s,H200 上从 266 降至 240 tokens/s)。在微观算子层面,HiSparse 使长上下文服务的整体吞吐量提升高达 3 倍;Flash Compressor 达到了 H200 峰值显存带宽的 80%,比朴素 PyTorch 实现快 10 倍以上;Lightning TopK 将小批次 Top-K 延迟压缩至约 15 微秒。在强化学习端,使用 32 张 GB300 跑 285B 模型(FP8 rollout + BF16 训练),首步 rollout 与训练的 log-prob 漂移仅约 0.023,奖励和评估分数稳步增长。该方案的代价与局限在于:RL 端的索引器重放机制目前仍是实验性功能,尚未经过端到端验证;此外,为了保证 RL 训练的绝对确定性(如固定 cuBLAS 工作区、关闭 TransformerEngine 非确定路径),系统牺牲了约 10% 到 15% 的训练吞吐量。对比基线在长上下文(200K+)下出现了超时失效,且其 MTP 投机解码在 H200 上因内核断言失败只能退化为 MTP-1。
一、 主要贡献
SGLang 和 Miles 团队在 DeepSeek-V4 发布首日(Day-0)即实现了对其推理和强化学习(RL)训练的全面支持。SGLang 和 Miles 构成了首个在发布当天服务和训练 DeepSeek-V4 的开源技术栈,其系统专门针对该模型的混合稀疏注意力架构、流形约束超连接(mHC)以及 FP4 专家权重进行了定制化构建。该技术栈全面支持了 DeepSeek-V4 的 1.6T Pro 版本和 284B Flash 版本。
在推理层面(缓存与注意力),系统引入了 ShadowRadix 前缀缓存、HiSparse CPU扩展 KV、带有图内元数据的 MTP 投机解码、Flash Compressor、Lightning TopK 以及分层多流重叠技术。在推理内核与部署层面,实现了快速的内核集成(包括 FlashMLA、FlashInfer TRTLLM-Gen MoE、DeepGEMM Mega MoE、TileLang mHC),支持 DP/TP/CP 注意力、DeepEP 上的 EP MoE 以及 PD 分离。在 RL 训练层面,实现了全并行策略(DP/TP/SP/EP/PP/CP)、Tilelang 注意力、增强的稳定性控制以及 FP8 训练。该系统广泛支持 Hopper、Blackwell、Grace Blackwell、AMD 以及 NPU 等硬件平台。
图1:在截取自《红楼梦》的 30K token 提示词下,SGLang 与其他开源引擎的解码吞吐量对比。针对每个引擎,均基于其官方建议尝试了尽力而为的投机解码配置。
二、 背景知识与模型关键特性
DeepSeek-V4(1.6T Pro, 284B Flash)在三个核心维度上对其前代 DeepSeek-V3.2 进行了扩展,这些架构特性直接驱动了底层系统的设计:
混合稀疏注意力(Hybrid sparse-attention):模型的每一层都将滑动窗口注意力(SWA,覆盖最后128个原始token)与两种压缩机制之一进行混合。这两种压缩机制分别为 C4(在 4:1 压缩的 KV 上进行 top-512 的稀疏注意力计算)或 C128(在 128:1 压缩的 KV 上进行密集注意力计算)。这种设计使得 1M token 的上下文窗口变得易于处理。
流形约束超连接(mHC, Manifold-Constrained Hyper-Connections):这是对标准残差连接的一种泛化设计,旨在改善梯度流并提升表示质量。
FP4专家权重(FP4 expert weights):模型原生采用 FP4 MoE 专家权重,以便在最新的 Blackwell 硬件上实现高效的模型服务。
三、 方法细节与系统设计
ShadowRadix:面向混合注意力的原生前缀缓存
混合注意力的缓存挑战:DeepSeek-V4 的每一层都将 SWA(针对最后128个原始token)与 C4(针对 4:1 压缩 KV 的 top-512 稀疏注意力)或 C128(针对 128:1 压缩 KV 的密集注意力)结合使用。此外,为了维护进行中的压缩 KV 槽,每个压缩层都设有一个状态池,用于存储正在进行中的压缩状态。这种复杂的机制打破了传统前缀缓存的假设:三个异构的 KV 池和两个压缩状态池必须在 Prefill(预填充)、Decode(解码)和投机解码过程中保持一致。图1展示了 $N = 1024$ 示例的逐层混合注意力作用域。
核心理念与基数树索引:为了解决上述一致性问题,系统引入了 ShadowRadix——一种专为混合注意力设计的原生前缀缓存机制。其核心理念是使用一棵基数树(radix tree)来索引虚拟全 token 槽(virtual full-token slots),这是一个由所有层共享的统一坐标系统。从每个槽中,系统将阴影(shadows,即每个池的索引映射)投影到物理池(SWA / C4 / C128)中。压缩状态的环形缓冲区(ring buffers)位于它们自己的池中,但第二级算术阴影会将每个环槽从 SWA 页面索引中映射出来——在逻辑上嵌套于 SWA 内部,但在物理上是独立的。这种设计使得生命周期得以解耦:一旦滑动窗口移过某个节点,墓碑机制(tombstoning)就会释放该节点的 SWA 槽,而其 C4/C128 的阴影则保持活跃且可共享。因此,一个 10k token 的请求只需保留 128 个 SWA token 加上其完整的 C4/C128 压缩 KV,而这些压缩 KV 正是其他前缀匹配请求可以重用的内容。
双计数器锁与安全匹配:由于阴影仅仅是同一源在不同池中的投影,每个池可以独立管理自己的生命周期。系统通过为每个节点引入双计数器锁(two-counter lock)来利用这一点:full_lock_ref 覆盖源(进而覆盖其 C4/C128 阴影),而 swa_lock_ref 仅跟踪该节点是否仍然落在某人的滑动窗口内。当 SWA 计数器归零时,系统会丢弃该节点的 SWA 槽,但该节点仍保留在树中,其压缩阴影继续为前缀匹配提供服务。匹配过程通过要求从匹配点开始有 128 个连续的活跃 token 延伸入窗口,来保持SWA安全(SWA-safe)。嵌套的环形阴影免费继承了这一特性:由于环槽的地址是 $swa\_page \times ring\_size + pos \pmod{ring\_size}$,释放一个 SWA 页面会自动使其内部的环失效,无需任何额外的跟踪。同时,环大小的选择确保了 MTP 回滚不会与环绕(wrap)发生竞争。
投机解码的兼容性:投机解码无需对该设计进行任何更改即可运行。唯一的问题是,草稿 token 在验证决定接受哪些 token 之前就被写入了环中,因此一个被拒绝后重试的步骤可能会发生环绕,并覆盖仍然属于活动窗口的槽。通过在投机下加倍环大小(C4: 8 -> 16, C128: 128 -> 256),系统留出了足够的余量,使得任何进行中的投机写入都会落在活动窗口之外,因此 EAGLE 实现了开箱即用。
投机解码(Speculative Decoding)
单层 MTP 头支持与图内融合:DeepSeek-V4 搭载了一个单层 MTP 头——这是一个独立训练的 DSv4 解码器层,仅运行 SWA 注意力(无压缩器,无索引器),并结合了上一步的隐藏状态($h\_proj$)和下一个 token 的嵌入($e\_proj$)作为其输入。系统在首日即支持了该功能;真正的系统级工作发生在更底层。混合注意力的元数据非常庞大,在调度器流上急切地准备这些数据会成为投机解码下的启动瓶颈。因此,系统将这种准备工作融合到了草稿(draft)和验证(verify)阶段的 CUDA 图(CUDA graph)中。
双重优化驱动加速:
- 图内元数据准备(In-graph metadata preparation):混合注意力在每个 pass 的元数据非常繁重——包括 SWA 页面索引、阴影映射的池槽、压缩器/索引器计划、每个池的写入位置等——但这些全部是基于页表和长度的索引算术,能够干净利落地适配设备端内核。因此,每个捕获的图只需要原始的批处理状态作为每次重放的输入(活跃请求、当前长度、新 KV 目的地),并复制到固定缓冲区中;捕获的内核在图内部重建其余部分,而 Python 在重放期间永远不会接触 per-pass 路径。这消除了原本会主导投机解码的每步启动开销。
- 重叠调度(Overlap scheduling):CPU 端的工作(结果处理、批次准备、内存释放)与 GPU 执行并行运行。
图3:混合稀疏注意力结合 ShadowRadix 和图内投机元数据,使得 SGLang 的解码吞吐量从 4K 到 900K 基本上保持平稳——接近模型完整的 1M 上下文窗口。在 B200(199 -> 180 token/s)和 H200(266 -> 240 token/s)上,性能下降均低于 10%。
HiSparse:使用分层内存为稀疏注意力提供涡轮增压
非活跃 KV 缓存的 CPU 卸载:近期引入的 HiSparse 是一种将非活跃 KV 缓存卸载到 CPU 内存的技术,从而为稀疏注意力实现更大的批大小和更高的吞吐量。HiSparse 与 C4 层天然契合:在每一步中,索引器的 top-k 仅接触一小部分压缩位置,因此绝大多数 C4 KV 在任何时刻都是非活跃的,可以存放在 CPU 上。由于 C128 是密集的(每个位置都会被接触)且 SWA 本身就很小(128 个 token),因此两者都无法从卸载中获益。通过使用 CPU 内存池来专门扩展 C4 KV 缓存池,系统将长上下文服务的整体 token 容量和吞吐量提高了最高 3 倍。
左图:GPU 仅保留小型设备缓冲区用于 C4 KV 缓存的活跃工作集,而一个更大的锁定 CPU 镜像则存储全上下文 KV 缓存。在每一步中,HiSparse 协调器从 CPU 换入未命中的页面,并使用 LRU 策略驱逐非活跃的 GPU 页面。新生成的 token 会异步备份到 CPU 镜像中。右图:在 2xB200 上,200K输入 / 20K输出,$swa\_full\_tokens\_ratio=0.001$ 配置下 DeepSeek-V4-Flash 的峰值吞吐量。
快速内核集成(Fast Kernel Integrations)
- 用于混合注意力的全新 FlashMLA 路径:DeepSeek-V4 的混合注意力将局部窗口上的 SWA 与 C4(在压缩 KV 上的稀疏 top-k)或 C128(在压缩 KV 上的密集计算)上的“额外”注意力结合在一起。系统集成了更新后的 FlashMLA 接口,使得 SWA 和额外注意力在一次融合的内核调用中运行(内核同时接收
k_cache和extra_k_cache及其各自的索引),并在前向传递中共享元数据构建。目标 GPU 为支持相应内核的 Hopper/SM90 和 Blackwell/SM100。 - 用于 MXFP8 x MXFP4 的 FlashInfer TRTLLM-Gen 融合 MoE:DeepSeek-V4 搭载了原生 FP4 专家权重,这使得小 batch 的 MoE 解码对专家权重带宽非常敏感。系统通过 FlashInfer 集成了 TRTLLM-Gen 融合 MoE 后端,以将 MXFP8 激活与 MXFP4 专家权重配对,并调整 SGLang 的权重/缩放布局以适应内核预期的格式。该路径针对 Blackwell GPU,并依赖于 Blackwell 特定的 FP4 张量核心机制以及分块/持久化执行。
- 带有 split-K 的 TileLang mHC 内核:mHC 用 $hc\_mult$ 个并行分支上的 per-token 混合替换了残差流。对 GEMM 输出进行 Sinkhorn 归一化以产生混合权重。在低延迟解码中,前置 GEMM 的 TileLang 内核可能会成为瓶颈,因为 batch size 很小且并行度受限。因此,系统扩展了两阶段 split-K TileLang 内核(
mhc_pre_gemm_sqrsum_splitk_kernel),将 K 维度划分到不同的 CTA 中,从而提高了小 batch size 下的 GPU 利用率。系统还集成了一个融合的mhc_pre_big_fuse_tilelang路径,该路径在启用 PDL 的情况下,将 RMSNorm、Sinkhorn 和残差混合结合在一个内核中。 - DeepGEMM Mega MoE 集成:系统在 Blackwell 上集成了 DeepGEMM 的 Mega MoE 路径。该内核将 EP 分发、第一个 FP8xFP4 专家 GEMM、SwiGLU、第二个 FP8xFP4 专家 GEMM 以及 EP 合并融合到一个基于对称内存的巨型内核(mega-kernel)中,使 NVLink 通信与张量核心计算重叠。由于 Mega MoE 需要带有 UE8M0 比例因子的转换后 FP4 专家权重,系统调整了现有的 EP/MoE 激活路径以直接消耗转换后的布局。这避免了在 GPU 内存中保留两份专家张量的常驻副本,同时仍启用了 DeepEP 路径和低延迟 Mega MoE 路径。
各种内核优化(Various Kernel Optimizations)
系统详细阐述了以下两个具有代表性的内核优化方案:
图5:两个通用的内核级优化。左图:Flash Compressor 将朴素的 5 阶段压缩链融合为一次片上计算,将 HBM 往返次数从 5 次削减至 2 次。右图:Lightning TopK 用 8 节点集群的 radix-select 归约替换了全局排序——每个 CTA 构建一个局部的 10 位基数直方图,集群对直方图进行归约以挑选出一个精确容纳 K = 512 个候选者的阈值,然后每个 CTA 仅分散(scatter)其高于阈值的条目。
Flash Compressor:IO感知精确压缩:压缩注意力通过在每个 token 的分数上进行 softmax 加权平均,将一组 token 压缩为一个 KV。一个朴素的实现会访问 HBM 五次,并在非连续的维度上运行 softmax,因此压缩本身的成本往往高于其所馈送的压缩注意力。系统将整个管道融合为一个单一的片上通道(图5,左),其中 C4 的 softmax 保持在 warp 级别,而 C128 则采用单一的 CTA 级别归约。在 H200 上,这可以达到高达 80% 的峰值内存带宽,并且比朴素的 PyTorch 管道快10倍以上。
Lightning TopK:稀疏注意力依赖于索引器中的 top-k 选择步骤。在 4:1 压缩下,1M token 的上下文仍然会为每个请求产生 256K 个候选者。在这个规模下,即使在 batch size 为 1 的情况下,朴素的实现也可能需要超过 100 微秒,这超过了上游索引器 GEMM 和下游稀疏注意力内核的延迟。系统通过为长上下文解码提供自定义的基数选择内核(radix-select kernel)来解决这个问题。该内核不执行完整的全局排序,而是缩小候选集,并使用带有异步拷贝的 CUDA 集群启动来增加并行度,使内存移动与选择重叠。这避免了昂贵的全局排序和通过 HBM 进行的跨 CTA 同步,有效地将小 batch top-k 延迟降低到约 15 微秒,并在大 batch size 下保持高性能。
并行性与部署(Parallelism and Deployment)
- 用于长上下文预填充的上下文并行(CP):DSv4 大部分混合注意力路径(SWA、C4/C128 压缩器、索引器)无法沿注意力头进行分片,因此 TP 在长提示词上会遇到瓶颈。CP 在注意力 rank 之间轮询分配 token:每个 rank 拥有
1 / cp_size的序列,并且所有 per-token 元数据(SWA 页面索引和长度、C4/C128 位置、页表、索引器 top-k 长度)都在本地重新索引,而压缩器的输出写入位置保持全局,以便生成的 KV 仍然连续着陆。CP 专门针对预填充 NSA 路径,并且是使长上下文 TTFT 在此扩展的关键。 - 高 TP 下的 FlashMLA 头填充:FlashMLA 要求在 Blackwell 上 $num\_heads$ 必须是 128 的倍数(在 Hopper 上为 64),而高 TP 大小很容易违反这一要求。系统没有重写内核,而是在张量级别进行每次调用的填充:分配一个全宽的 Q 缓冲区,将 rank 的真实头复制到其切片中,填充其余部分,调用内核,然后裁剪输出。由于 MLA 解码通常在 KV 读取上受限于内存带宽,因此额外的 FLOPs 带来的开销可以忽略不计。
- 用于 PD 分离的分页 KV 传输:系统使用页面索引的 KV 传输扩展了 SGLang 现有的 PD 传输协议。页面通过标准路径按索引移动,并在接收端通过其自身的阴影映射重新解释,因此传输层对 DSv4 设备上的非均匀布局保持无感知(见 ShadowRadix)。
- DeepEP + Mega MoE 上的专家并行(EP):DSv4 的 MoE 运行在 DeepEP 之上以实现全对全的专家路由。在此之上,系统集成了 DeepGEMM 的 Mega MoE 路径。Mega MoE 期望专家权重采用与标准 DeepEP 路径不同的布局;系统将 Mega MoE 权重设置为 DeepEP 专家张量的别名视图,以便大规模部署不会在 GPU 内存中为两份专家权重付出代价。完整的路径将其与 EP 端的 SwiGLU-clamp 和 JIT 激活融合结合在一起。
分层多流重叠(Hierarchical Multi-Stream Overlap)
注意力准备阶段由一群小内核组成——Q/KV 投影、压缩器、索引器——在小 batch 时,每个内核都无法填满 GPU。此时,串行启动的开销开始主导解码过程。
系统在两个层面上将它们跨 CUDA 流扇出(图6)。顶层并行运行四个准备操作。索引器是其中最繁重的,因此在其内部再次进行拆分——Q 投影、权重投影和压缩器 GEMM 也相互重叠。诸如 $q\_lora\_ready$ 和 $q\_scale\_ready$ 等事件一旦完成依赖关系就会在流之间进行交接,因此没有任何操作需要等待整个流排空。
图6:用于注意力准备的两级流扇出,具有细粒度的 $q\_lora\_ready$ / $q\_scale\_ready$ 事件,跨流链接依赖的工作。
该机制在 CUDA 图内部运行,且仅在小 batch 时触发——因为大的预填充已经使设备饱和,重叠不会带来任何额外收益。
四、 补充细节(强化学习 —— Miles 支持)
图4:DeepSeek-V4 的首日验证 RL 管道:Miles 和 SGLang 共同提供了稳定性(R3 和索引器重放,Step-0 训练-推理差异约 0.02-0.03)、效率(DP/TP/SP/EP/PP/CP,Tilelang 内核,FP8/BF16 rollout 和训练)以及广泛的硬件支持(Hopper, Blackwell, Grace Blackwell)。
训练后端:Megatron-LM 中的 DeepSeek-V4 建模
对于 DeepSeek-V4,由于全新的模型架构(例如复杂的压缩注意力、新的稀疏-MLA 和索引器模块以及 mHC 层),大部分算子被重新构建。
支持 DP/TP/SP/EP/PP/CP 并行:支持 Megatron 中的所有六种策略。
- TP 和 SP 在压缩注意力模块中实现。
- PP:mHC 的四个流必须跨越阶段边界,以便接收阶段可以在每个子层混合它们,因此系统通过 p2p 传递 $[seq, batch, hc\_mult, hidden]$ 格式,而不是通常的 3-D 张量。
- CP:系统在 DeepSeek-V4 压缩注意力上支持 all-gather CP。C4 压缩器的重叠变换将每个压缩组的一半沿序列向后移动一个位置(相邻组在没有额外 KV 成本的情况下共享上下文),因此在 CP 下,这种移动跨越了 rank 边界。仅靠光环交换(halo exchange)就能修复重叠,但下游索引器 top-k 无论如何都需要每个压缩位置,因此系统将这两个集合操作折叠成一个单一的 all-gather。例如,在 CP=4 且序列为 1024 token 的情况下,每个 rank 持有 256 个 token;系统跨 CP 组 all-gather 完整的压缩张量,在组装后的序列上应用重叠,然后切片返回本地视图。C128 层没有重叠变换,完全跳过 all-gather。
内核支持:基于 Slime 为 GLM-5 提供的内核支持,系统针对 DeepSeek-V4 进行了采用和扩展。
- 调整 Tilelang DSA 索引器内核以适应 DeepSeek-V4 的索引器架构,并将其用于闪电索引器中。
- 扩展稀疏-MLA Tilelang 内核,在 softmax 分母中加入每头可学习的注意力-sink logits 以支持核心注意力。
- 使用融合的 Tilelang Sinkhorn 内核用于 mHC。
RL 特性
DeepSeek-V4 训练栈支持并测试了重要的 RL 特性,包括 Rollout Routing Replay (R3)、TIS/MIS、FP8 训练和 FP8 注意力 QAT:
- 支持 FP8 rollout,FP8/BF16 训练:支持 FP8 rollout 以及 BF16 和 FP8 训练,并在权重更新中包含量化处理器。
- 注意力 QAT:在 rollout 时以 FP8 服务的路径(压缩器 KV、普通 KV、索引器查询)上模拟 FP8 激活量化,以在使用 FP8 rollout 精度时与 SGLang 实现内核级数值匹配。
- Rollout Routing Replay (R3):在 miles 中将 R3 处理器和 TIS/MIS 损失扩展为 $(b, s, h, d)$ 格式,以支持 DeepSeek-V4 后端。
- 索引器重放(实验性):DeepSeek-V4 在 MoE 路由之外拥有第二种随机算子——DSA 索引器在压缩 KV 上的 top-k。系统实现了索引器重放机制,在压缩注意力后端捕获 SGLang 中的 rollout top-k,在现有 rollout 通道上传输它们,并在训练时每个 C4 层重新注入。这是一个实验性功能,管道通过了短上下文正确性检查,但尚未进行端到端验证。
混合精度堆栈上的数值精度
- 系统在训练、检查点转换和权重更新过程中仔细维护了敏感主权重及其梯度的 FP32 精度。
- 压缩器的反向传播是在长轴上进行 softmax 加权求和,其中最小的加数至关重要;跨 TP rank 累积的 BF16 舍入会使总和偏向幅度最大的贡献者。系统仅将此 all-reduce 切换为 FP32。
- 选择性地冻结不稳定的路径:mHC 中的 Sinkhorn(作为混合预言机)、MoE 路由器门、每专家分数校正偏差,以及哈希路由的早期层。
- 为了避免 KL 损失中出现随机数值尖峰,系统固定了一些确定性操作:cuDNN 确定性、NCCL 设置为 Ring、禁用 TransformerEngine 的非确定性路径,并将 cuBLAS 绑定到固定工作区。这消耗了大约 10–15% 的吞吐量。
Megatron 中的多项修复:
- 修复了混合 FP32/BF16 组下 distributed-optimizer 中的检查点精度问题。
- 为了解决恢复检查点时遇到的 OOM 问题,系统在加载期间在 CPU 上分配优化器状态,以避免出现巨大的峰值内存。
五、 实验环境
硬件配置:
- 推理基准测试:NVIDIA B200 平台(用于 1.6T Pro,TP=8);NVIDIA H200 平台(用于 285B Flash,TP=4)。
- RL 训练:32 张 NVIDIA GB300 GPU,配置 TP/SP/EP/PP 并行。
- 系统整体声明支持的硬件包括 Hopper, Blackwell, Grace Blackwell, AMD 以及 NPU。
软件配置:
- 核心代码库基于 SGLang 和 Miles,RL 训练后端基于 Megatron-LM 进行深度定制。
- 采用 FP8 rollout + BF16 训练精度,并启用 R3 机制。
数据集与任务规模:
- 推理测试:截取自《红楼梦》的文本,生成 30K token 的提示词;针对长上下文测试,最高测试到 900K token。
- RL训练:使用 DAPO 算法,最大响应长度设置为 4096。
六、 实验结果
推理基准测试(对比其他 OSS 引擎):
- 单 Batch 解码吞吐量(图1):在 30K token 上下文下,SGLang 在 B200 (Pro) 和 H200 (Flash) 上均展现出显著高于其他开源引擎的解码吞吐量(解码吞吐量计算为
1000 / TPOT (ms))。 - 投机解码表现:SGLang 采用 EAGLE 3/1/4 配置,在 Pro 和 Flash 模型上的接受长度均约为 2.5。其他 OSS 引擎在 B200 Pro 上采用 MTP-3,接受长度仅约 1.19(位置 0/1/2 接受的 token 数严重偏斜至 2226/354/55,表明其 MTP 路径在 Pro 上未达到完全有效性);在 H200 Flash 上,其他引擎因内核断言错误回退到 MTP-1,接受长度约 1.92。
- 长上下文稳定性(图3):SGLang 在 4K 到 900K 的上下文中保持了极高的性能稳定性,吞吐量下降幅度不足 10%。相比之下,其他开源引擎在 200K+ 输入时未能在此轮测试超时前返回结果。
HiSparse 性能提升(图6):
在 DeepSeek-V4-Flash(2xB200,200K输入,20K输出)的配置下,启用 HiSparse 后,峰值吞吐量相比基线提升了最高 3 倍。
RL 训练结果(图10、图11):
在 32 张 GB300 上启动 285B 模型训练,训练动态表现稳定:
- 在第一步时,rollout 与 training 的对数概率漂移(log-prob drift)极小,约为 0.023。
- 随着训练步数的增加,Rollout 原始奖励(raw reward)和 AIME 评估准确率均呈现出稳步增长的趋势。
七、 结论与展望
SGLang 和 Miles 团队成功在 DeepSeek-V4 发布首日交付了针对该模型复杂架构(混合注意力、mHC、FP4)高度优化的推理与 RL 训练全栈开源解决方案。通过引入 ShadowRadix、图内投机解码元数据、HiSparse 卸载以及多项底层内核重写与并行策略,系统不仅实现了极高的推理吞吐量与长上下文支持,还确保了 RL 训练的数值稳定性和高效性。
未来的工作路线图将继续在开源社区推进,相关进展可通过 SGLang (issue #23602) 和 Miles (issue #1046) 的 GitHub 仓库进行追踪。团队特别致谢了 DeepSeek 提供的模型及首日合作,RadixArk 提供的跨越 Hopper、Blackwell 及 Grace Blackwell 的 GPU 基础设施支持,以及 SGLang 和 Miles 社区的贡献者们。
引用信息:
【1】Ke Bao et al., "DeepSeek-V4 on Day 0: From Fast Inference to Verified RL with SGLang and Miles", 2026. URL: https://www.lmsys.org/blog/2026-04-25-deepseek-v4/
💬 评论讨论
欢迎在这里分享您的想法和见解!