DualPath: Breaking the Storage Bandwidth Bottleneck in Agentic LLM Inference

Yongtong Wu1,3 Shaoyuan Chen2,3 Yinmin Zhong1,3 Rilin Huang1 Yixuan Tan3 Wentao Zhang3 Liyue Zhang3 Shangyan Zhou3 Yuxuan Liu3 Shunfeng Zhou3 Mingxing Zhang2 Xin Jin1 Panpan Huang3
1School of Computer Science, Peking University 2Tsinghua University 3DeepSeek-AI

1. 主要贡献

本文针对代理式(Agentic)大语言模型(LLM)推理中存在的 I/O 瓶颈问题,提出了一种名为 DualPath 的推理系统。作者首先识别出在多轮对话的代理式工作负载中,KV-Cache 的 I/O 开销而非计算已成为主要性能瓶颈。在现有的预填充-解码(PD)分离架构中,仅由预填充引擎(Prefill Engines, PEs)从外部存储加载海量 KV-Cache 会导致存储带宽的不平衡:PE 端的存储网卡(SNIC)带宽饱和,而解码引擎(Decoding Engines, DEs)的存储网卡却处于空闲状态。

为了解决这一问题,DualPath 的核心创新点在于引入了双路径 KV-Cache 加载机制。除了传统的“存储到预填充”路径外,DualPath 启用了一条新的“存储到解码”路径,即允许将 KV-Cache 加载到解码引擎中,随后通过计算网络上的 RDMA 高效传输给预填充引擎。这种设计利用了全局存储带宽池,打破了单侧瓶颈。

图 1:现有瓶颈(左)和 DualPath(右)。
图 1:现有瓶颈(左)和 DualPath(右)。

DualPath 系统结合了以下关键设计:
1. 优化的数据路径:能够动态选择加载路径,避免网络拥塞。
2. 以 NIC 为中心的流量管理:隔离 KV-Cache 流量与延迟敏感的模型执行通信。
3. 全局调度器:在预填充和解码引擎之间动态平衡负载。

实验结果表明,在具有真实代理工作负载的三种模型评估中,DualPath 在内部推理系统上的离线推理吞吐量提高了 1.87 倍,在不违反服务等级目标(SLO)的情况下,在线服务吞吐量平均提高了 1.96 倍

2. 背景知识与关键观察

LLM 推理基础
LLM 推理正从单轮对话演变为多轮交互的代理系统。现代推理系统普遍采用 PD 分离架构 [37, 58],将预填充阶段(计算密集、批处理)与解码阶段(内存受限、延迟敏感)分配给专用的 PE 和 DE。为了克服 HBM 容量限制并提高 GPU 利用率,系统常采用 分层预填充(Layerwise Prefill) [17, 50] 技术,即按层分配和释放 KV-Cache。

图 2:代理轨迹示例。
图 2:代理轨迹示例。

代理式 LLM 的使用模式
如图 2 所示,代理式应用涉及与环境的多轮交互。其典型特征是长上下文、短追加(Short-append)和多轮次。由于上下文在各轮次间累积,且 >95% 的 token 通常是重复的(KV-Cache 命中),每次只需对新追加的 token 进行预填充计算。这种模式导致了极高的 KV-Cache 命中率,使得 KV-Cache 加载效率取代纯计算成为主导性能因素。由于上下文极长(可能达到百万级 token),必须依赖容量大但速度较慢的外部 SSD 存储 [13]。

现代 AI 数据中心架构
典型的 AI 集群节点(如 NVIDIA DGX SuperPOD)配备专用的计算网卡(CNIC,用于 GPU 间互联)和存储网卡(SNIC,用于访问数据和模型)。计算网络与存储网络物理隔离 [55] 是基本原则,旨在防止干扰并最大化性能。

瓶颈分析与动机
通过调查,作者发现代理式推理任务中 GPU 利用率严重不足,主要由三个因素导致:
1. 极高的 I/O 需求:代理工作负载的高 KV-Cache 命中率导致缓存-计算比(cache-compute ratio)极高。例如,DeepSeek-V3.2 的该比率约为 22 GB/PFLOP,远超硬件能力,造成存储带宽瓶颈(见 Table 1)。
2. 硬件演进趋势不匹配:如图 3 所示,从 Ampere 到 Blackwell,I/O-计算比下降了 14.4 倍。网络带宽和 HBM 容量的增长滞后于 GPU FLOPS 的增长,导致在代理工作负载下撞上内存墙和通信墙。
3. 存储网络利用率不平衡:在现有 PD 分离系统中,命中 token 的 KV-Cache 仅由 PE 直接从存储加载。这导致 PE 端的 SNIC 持续饱和,而 DE 端的 SNIC 大量闲置,无法利用集群的聚合存储带宽。

表 1:在上下文长度(16k–64k)下,追加长度为 429 时的缓存-计算比。除非另有说明,KV-Cache 数据类型默认为 FP8。
表 1:在上下文长度(16k–64k)下,追加长度为 429 时的缓存-计算比。除非另有说明,KV-Cache 数据类型默认为 FP8。
图 3:左:NVIDIA GPU 的硬件趋势。右:随着请求批大小变化(每个请求有 30K 上下文和 300 个追加 token),相对 token 吞吐量的变化。
图 3:左:NVIDIA GPU 的硬件趋势。右:随着请求批大小变化(每个请求有 30K 上下文和 300 个追加 token),相对 token 吞吐量的变化。

基于上述分析,利用 DE 闲置的 SNIC 带宽加载 KV-Cache,并通过具有大量空闲带宽的计算网络回传给 PE,成为打破瓶颈的关键机会。

3. 方法细节

DualPath 系统基于 PD 分离和分层预填充技术构建。系统包含推理引擎(PEs 和 DEs)、流量管理器(负责数据传输)和请求调度器(负责任务分配和路径选择)。

3.1 双路径加载机制

DualPath 引入了双路径加载架构,动态利用所有引擎的存储带宽。为实现此机制,系统在每个 PE 和 DE 上分配少量的 DRAM 作为缓冲区(PE buffer 和 DE buffer)。

图 4:双路径加载示意图。调度器动态在两条路径之间分配数据流量。
图 4:双路径加载示意图。调度器动态在两条路径之间分配数据流量。

预填充阶段的 PE 读取路径(PE read path)
如 Fig 4a 所示,该路径遵循传统模式但增加了中转步骤:
1. 存储读取:命中 token 的 KV-Cache 从持久化存储读入 PE buffer(Label 1-2)。
2. HBM 加载与计算:在计算某一层 Attention 之前,该层的 KV-Cache 流式传输至 PE HBM(Label 3-4),用于计算 prompt 中未命中(miss)token 的 KV-Cache。
3. 传输至 DE:所有 KV-Cache(命中和未命中部分)随后传输至 DE buffer,以构建完整的 prompt KV-Cache(Label 5-7)。
4. 过程循环:步骤 3-7 会重复 $N_{layers}$ 次。在预填充前向传递过程中,传输与计算重叠。

预填充阶段的 DE 读取路径(DE read path)
如 Fig 4b 所示,这是 DualPath 新增的核心路径:
1. 存储读取:命中 token 的 KV-Cache 首先读入 DE buffer(Label 1-2)。
2. 反向传输:在 PE 进行预填充时,对应层的 KV-Cache 从 DE buffer 读取并传输给 PE(Label 3-5),该过程与计算重叠。
3. 过程循环:此过程重复 $N_{layers}$ 次。
4. 合并:层计算完成后,仅未命中 token 的 KV-Cache 传输至 DE buffer 并与现有的命中 token KV-Cache 合并。

解码阶段(Decode Phase)
当 DE buffer 接收到完整的 prompt KV-Cache 后,解码阶段开始。DE 分配 HBM 并执行主机到设备(H2D)的传输(Fig 4a Label 8-9; Fig 4b Label 6-7),随后释放 CPU 内存。这种设计虽然增加了 DRAM 和 CNIC 的带宽压力,但 DE buffer 的存在有效减少了 GPU 显存占用,这对处理长上下文至关重要。解码过程中,每累积一个完整的 token 块(如 64 tokens),会立即持久化到磁盘。

KV-Cache 块布局
系统采用两种布局:Full Block(包含所有层)和 Layer Block(包含单层)。与存储的所有交互均采用 Full Blocks。而在 PE 读取路径和 DE 读取路径的流式传输中,均使用 Layer Blocks。

3.2 无瓶颈分析

作者通过理论推导证明了 DualPath 在合理的 P/D(预填充/解码)比例下不会引入计算网卡(CNIC)或 DRAM 瓶颈。

PE CNIC 带宽分析
对于 PE CNIC,由于存在回环流量(loopback traffic),重点分析 PCIe 侧压力。
读取操作的总流量为:
$2 \times T_p \times Dg = 2Bs/g \le B$
由于存储带宽 $s$ 总是小于计算网络带宽 $g$,读取方向无瓶颈。
写入操作的总流量为:
$(T_p + T_c) \times Dg = Bs/g \times (1 + D/P) \le B$
推导出无瓶颈条件为:

$P/D \ge \frac{s}{g-s}$

DE CNIC 带宽分析
对于 DE CNIC,读取操作流量为:
$(T_p + T_c \times 2) \times Pg = s/g \times (P/D + 2) \times B \le B$
推导出:

$P/D \le \frac{g-2s}{s}$
写入操作流量为:
$(2T_p + T_c) \times Pg \le B$
推导出:

$P/D \le \frac{g-s}{2s}$

DRAM 压力分析
综合 PE 和 DE 的 DRAM 读写压力(半双工),要求 DE MEM 压力 $\le M$(内存带宽),得出:

$P/D \le \frac{M/Bs - 3}{2}$

总结
综合上述分析,无瓶颈的 P/D 比例范围为:

$\frac{s}{g-s} \le P/D \le \min \left\{ \frac{g-2s}{s}, \frac{g-s}{2s}, \frac{M/Bs - 3}{2} \right\}$
在实际配置中(如 $g \approx 500$ GB/s, $s/N_{gpu} \approx 50$ GB/s),无瓶颈范围约为 $1 \le P/D \le 7$,覆盖了绝大多数实际场景。

3.3 以 CNIC 为中心的流量管理器

为了解决 KV-Cache 传输可能干扰延迟敏感的模型执行通信(如 Expert Parallel 的 AllToAll)的问题,DualPath 采用了 以 CNIC 为中心(CNIC-centric) 的数据传输方案。

流量隔离
系统利用 InfiniBand 的 虚拟通道(Virtual Lanes, VLs) 实现隔离。
1. 高优先级 VL:分配给所有模型推理通信流量。
2. 低优先级 VL:分配给 KV-Cache 传输等其他流量。
3. 配置:交换机和网卡的 VL 仲裁器配置为加权轮询(Weighted Round Robin),保留约 99% 的带宽给高优先级 VL,剩余带宽给低优先级 VL 以防饥饿。这种配置确保了模型流量不受干扰,同时 KV-Cache 流量能利用空闲带宽。该设计同样适用于 RoCE 网络(利用 DSCP 和队列)。

CNIC 辅助的 KV-Cache 拷贝
现有的 GPUDirect Storage 或 CUDA copy engine 无法有效隔离 KV-Cache 流量与高优先级通信。DualPath 强制所有进出 GPU 的数据流量(包括本地 H2D/D2H)都经过配对的 CNIC。
1. 读取流程:KV-Cache 先从存储读入 Host DRAM,然后提交 RDMA Write 请求给本地 CNIC 执行 H2D 拷贝(Loopback)。
2. 写入流程:反之亦然。
3. 优势:确立 CNIC 为所有 GPU PCIe 流量的中央 QoS 调度器,利用其 VL 仲裁器优先处理推理流量。虽然路径看似绕弯,但测量显示提交 RDMA Write 请求(约 1 $\mu s$)比 cudaMemcpyAsync(5-7 $\mu s$)开销更低,且支持 Doorbell Batching。

3.4 自适应请求调度器

为了同时平衡 NIC 流量和 GPU 利用率,调度分为两级:引擎间调度(Inter-Engine)和引擎内调度(Intra-Engine)。

引擎间调度(Inter-Engine Scheduling)
引擎被组织成组,仅由 Leader Engine 与调度器交互。
1. PE 调度
* 指标报告:引擎报告未完成请求数 $N_{req}$、总 token 数 $N_{tok}$ 和节点磁盘读取队列长度 $L_{disk\_read}$。
* 分类:将 PE 分为三类:(1) 过载 ($N_{tok} > \beta$);(2) 短读取队列节点 ($L_{disk\_read} \le \alpha$);(3) 长读取队列节点。
* 策略:优先分配给类别 (2) 中 $N_{tok}$ 最小的 PE,其次是类别 (3)。不分配给过载 PE。这有助于平衡存储 NIC 的负载。逻辑见 Fig 5。

图 5:引擎间 PE 调度示意图。所有八个 GPU 都在同一个 PE 引擎组中,调度器将选择最佳的一个。
图 5:引擎间 PE 调度示意图。所有八个 GPU 都在同一个 PE 引擎组中,调度器将选择最佳的一个。
  1. DE 调度

    • 阶段 1(跨组):将请求分配给总 $N_{tok}$ 最小的 DE 组,平衡组间负载。
    • 阶段 2(组内):计算组内 DE 的剩余 HBM。定义高 token 阈值 $H$。优先选择剩余 HBM 足够且 $N_{tok}$ 较小的 DE,同时尽量避免选择已处于高负载(Category 1)的 DE,以防止抢占风险。
  2. 路径选择:在确定 PE 和 DE 后,选择读取队列较短的一侧进行 KV-Cache 读取。

引擎内调度(Intra-Engine Scheduling)
仅 PE 需要此调度。由于数据并行(Data Parallelism)可能导致不同 GPU 处理的请求长度不一,进而导致 Attention 层执行时间不同步,产生气泡(Bubbles)。

图 6:引擎内调度。左:基于计算配额的批次选择。右:应用计算配额前后的 GPU 时间线。
图 6:引擎内调度。左:基于计算配额的批次选择。右:应用计算配额前后的 GPU 时间线。
  1. 层时间预估:使用 $(N_{cached}, N_{new})$ 对描述请求,预估 Attention 层执行时间。
  2. 计算配额(Compute Quota)算法:设定计算时间上限。按 FIFO 顺序加入请求,直到预计时间超过配额。
  3. 分块预填充(Chunked Prefill):如果加入某请求会导致超时,则对该请求的 $N_{new}$ 进行二分查找,仅处理一部分 token 以填满剩余配额(Fig 6)。这最大限度地减少了并行的同步等待时间。

4. 实验环境

  • 数据集

    • 从生产环境的代理式 RL 训练工作负载中收集了三个数据集(32K, 48K, 64K MaxLen)。
    • 特点:平均交互 60-157 轮,平均追加长度 429-608 tokens,KV-Cache 命中率极高(见 Table 2)。
  • 模型

    • DeepSeek V3.2 660B (MoE, Sparse Attention).
    • DeepSeek 27B (Downscaled version).
    • Qwen2.5-32B (Dense, GQA).
  • 硬件配置

    • NVIDIA Hopper GPU(每节点 8 卡)。
    • 网络:每 GPU 配备 1 个 400Gbps RDMA CNIC(物理隔离的计算网络);每节点配备 1 个 400Gbps SNIC(存储网络)。
    • 存储:Cluster-wide 3FS 分布式存储,无内部 DRAM 缓存,可跑满 SNIC 带宽。
  • 软件配置

    • 基于内部推理框架(集成 FlashMLA, DeepGEMM, DeepEP)。
    • 对比基线:
      • SGL(MC): SGLang + HiCache + Mooncake Store。
      • Basic: 未修改的内部框架(仅使用 PE 侧加载)。
      • Oracle: DualPath 的理论上限版本(旁路所有磁盘读取和传输)。
表 2:代理轨迹数据集统计。
表 2:代理轨迹数据集统计。

5. 实验结果

离线批量推理(Offline Batch Inference)
该场景模拟 RL 训练中的 Rollout 阶段。

  • 吞吐量提升:如图 7 所示,DualPath 在 DeepSeek 660B 上相较于 Basic 实现了最高 1.87 倍 的加速,且性能接近 Oracle,表明 KV-Cache I/O 开销被极大消除。在 DS 27B 和 Qwen 32B 上也有显著提升(最高 1.78x)。
  • 参数影响:随着 Agent Batch Size 和最大上下文长度(MAL)增加,DualPath 的优势更明显。
  • 追加/生成长度影响:DualPath 在短追加和短生成场景下优势最大(Fig 9)。随着追加长度增加,计算占比提升,I/O 瓶颈减弱,Basic 与 DualPath 差距缩小。
  • P/D 比例影响:在 DS 27B 的测试中(Fig 8),DualPath 在所有 P/D 比例下均优于 Basic。DualPath 1P1D 的性能与 Basic 2P1D 相当,证明了 DualPath 有效利用了全集群存储带宽。
图 7:不同代理数量和最大代理上下文长度下的离线推理性能。上:DS 27B。中:DS 660B。下:Qwen 32B。N/A 表示在完成前遇到错误。
图 7:不同代理数量和最大代理上下文长度下的离线推理性能。上:DS 27B。中:DS 660B。下:Qwen 32B。N/A 表示在完成前遇到错误。
图 8:预填充-解码比例对离线推理性能的影响(DS 27B)。
图 8:预填充-解码比例对离线推理性能的影响(DS 27B)。
图 9:左:不同追加长度(DS 660B,64K 上下文,1024 个代理)。右:不同生成长度(DS 660B,64K,1024 个代理)。
图 9:左:不同追加长度(DS 660B,64K 上下文,1024 个代理)。右:不同生成长度(DS 660B,64K,1024 个代理)。

在线服务(Online Serving)
该场景评估系统在满足 SLO(TTFT $\le$ 4s, TPOT $\le$ 50ms)下的服务能力。
* 容量提升:如图 10 所示,DualPath 支持的每秒代理到达率(APS)显著高于 Basic(DS 660B 提升 2.25 倍)。
* 延迟分析:DualPath 的 TTST 与 Basic 相当,TPOT 也没有因为引入复杂的 DE 路径而增加,证明了流量隔离的有效性。TTFT 的分解显示(Fig 12 左),Basic 的排队时间因存储带宽不足而剧增,而 DualPath 保持稳定。
* 平均完成时间:DualPath 显著降低了平均任务完成时间(JCT)(Fig 11)。

图 10:TTFT、TTST 和 TPOT 随代理到达率(APS)变化的函数。阴影表示实验结束前最后 150 秒的波动。上:DS 27B,下:DS 660B。
图 10:TTFT、TTST 和 TPOT 随代理到达率(APS)变化的函数。阴影表示实验结束前最后 150 秒的波动。上:DS 27B,下:DS 660B。
图 11:在线服务的所有轨迹的平均完成时间与到达率的关系。
图 11:在线服务的所有轨迹的平均完成时间与到达率的关系。
图 12:左(§7.4):在线服务(DS 660B)在不同 APS 下的 TTFT 分解。Sch. 为调度,A. 为分配,R. 为读取 KV-cache,PF. 为预填充。每对柱状图中,第一个为 DualPath,第二个为 Basic。右(§7.5):离线推理消融结果(DS 660B,64K 上下文长度)。Layer, DPL, Sched 分别代表分层预填充、DualPath 加载和调度。
图 12:左(§7.4):在线服务(DS 660B)在不同 APS 下的 TTFT 分解。Sch. 为调度,A. 为分配,R. 为读取 KV-cache,PF. 为预填充。每对柱状图中,第一个为 DualPath,第二个为 Basic。右(§7.5):离线推理消融结果(DS 660B,64K 上下文长度)。Layer, DPL, Sched 分别代表分层预填充、DualPath 加载和调度。

消融研究
通过逐步添加技术组件(Fig 12 右):
1. +Layerwise Prefill:减少 JCT 17.21%(缓解 HBM 瓶颈)。
2. +DualPath Loading:进一步大幅减少 JCT 38.19%(利用分布式存储带宽)。
3. +Scheduling:最终减少 JCT 45.62%(实现负载均衡)。

负载均衡
调度算法将存储 NIC 的流量不平衡度从 1.53 降低到 1.18(Fig 13)。Attention 层执行时间的最大/平均比率保持在 1.06 的低位(Fig 14),有效减少了计算气泡。

图 13:存储 NIC 流量的负载均衡情况
图 13:存储 NIC 流量的负载均衡情况
图 14:Attention 执行时间的负载均衡情况。
图 14:Attention 执行时间的负载均衡情况。

大规模扩展性
在 1152 个 GPU 的规模下(Table 3),DualPath 保持了线性加速比和稳定的延迟,证明了其在大规模集群上的可用性。

表 3:大规模实验结果。
表 3:大规模实验结果。
图 15:48P96D 离线推理指标。1e7 是 Prompt TPS 的缩放因子。
图 15:48P96D 离线推理指标。1e7 是 Prompt TPS 的缩放因子。

6. 结论

本文提出的 DualPath 框架成功解决了 PD 分离架构下代理式 LLM 推理的 KV-Cache 读取不平衡问题。通过创新的双路径加载机制、以 NIC 为中心的流量隔离和自适应调度,DualPath 有效聚合了集群的存储带宽,打破了预填充端的 I/O 瓶颈。实验证明其在离线推理和在线服务场景下均带来了近 2 倍的性能提升,为高效的代理式 AI 系统提供了重要的架构参考。

7. 附录

流量隔离配置细节
* InfiniBand:配置 4 个虚拟通道(VL)。High-limit 设为 240,确保高优先级流量(VL0, VL1, VL3)在加权轮询中占据主导地位,低优先级(VL2)分配少量带宽防止饥饿。
* RoCE:使用 DSCP 映射到硬件 Traffic Classes (TC),配置 4 个无损 RDMA TC 并启用 PFC,带宽分配逻辑与 IB 类似。

模型规格
DeepSeek 27B 模型的详细参数:Hidden dim 2560,30 层,32 Heads,72 个路由专家(每 token 激活 6 个)。

KV-Cache 块布局
采用 Layer Block(单层 [1, tokens, size])和 Full Block(全层 [layers, tokens, size])设计。存储交互使用 Full Block,传输使用 Layer Block,通过简单的拼接操作避免内存布局转换开销。存储结构采用 Trie 树。