cluster-llm-deployment

Cross-category topic | 84 sources

Cluster LLM 部署:存储、计算、执行与成本 #

§1 主题缘起 #

大规模 LLM 集群部署正在经历从"单节点单模型"到"千 GPU 级多模态流水线"的范式跃迁。这一跃迁受到六股力量的同时驱动。

存储维度:vLLM 的 PagedAttention 将 OS 虚拟内存分页引入 KV cache 管理,以 block table + CoW 消除碎片实现 2–4× 吞吐提升 [2309.06180],奠定了几乎所有后续系统的内存管理基础。Mooncake 进一步将 KV-Cache 从 GPU 扩展到分布式 CPU DRAM/SSD pool,实现跨节点 prefix 复用和 PD 分离,长上下文吞吐提升 525% [2407.00079]。Agentic workload 的兴起使 KV-Cache 命中率极高(DeepSeek 生产 coding agent trace 显示 98.7% 命中率 [2602.21548]),prefill 从计算密集退化为 I/O 密集。PD 分离架构将所有存储读取集中在 prefill engine 的 SNIC 上,使存储带宽成为单点瓶颈 [2602.21548]。DeepSeek-V2 的 MLA(Multi-head Latent Attention)将 KVCache 压缩 93.3%——仅缓存 512 维 latent 而非完整 K/V [2405.04434],这一创新使新一代 hybrid attention 模型(KDA:MLA 混合)将单实例 KV 吞吐从 60 Gbps 压缩到 3 Gbps,跨数据中心 PD 分离从不可行变为可行 [2604.15039]。KV-Cache 传输瓶颈催生了多路径解决方案——KVServe 用 service-aware 压缩实现最高 10× KV 数据削减 [kvserve],PPD 发现 Turn 2+ 的 append-prefill 仅造成 2% TPOT 劣化(vs full prefill 的 48%),可安全回路到 decode 节点 [2603.13358]。ForkKV 则发现 multi-LoRA agent 场景下 KV cache 因 adapter 差异而无法共享——通过 DualRadixTree + CoW 将 base/residual KV 分离,实现 12.7× 单 agent 内存缩减和 1.25–3.04× 吞吐提升 [2604.06370]。RTP-LLM 将 KV cache 管理扩展为四层层级架构——GPU 显存 → 本地 CPU 内存 → 远程 CPU(RDMA)→ 分布式存储(3FS),统一 hash map 实现 $O(B)$ 前缀匹配替代 $O(B \times W)$ 逐 worker 查找,生产环境下 cache 命中率达到 45%(vs vLLM 19%、SGLang 29%),75% prefill 机器缩减 [2605.29639] [2605.29639]。RL 训练中的权重传输同样面临 TB 级数据搬运挑战——Ray object store 传 40 GB 需要 32 秒,而 GPU-direct RDMA 只需 0.2 秒 [2604.09107]。KV cache 管理这一子领域已成熟到催生系统性综述——KV Cache Management Survey 将全部工作归纳为 token-level / model-level / system-level 三层分类法,为该方向建立了统一的问题框架 [2412.19442]。传输通道之外,KV cache 的编码与流式加载本身也是一条独立瓶颈:CacheGen 将 KV cache 视作可量化编码的张量流,用算术编码 + 分层量化把 KV 表示压缩到 3.5–4.3× 更小并支持带宽自适应流式加载,TTFT 相比重计算降低 3.2–3.7× [2310.07240]。LMCache 则把"KV cache 层"抽象为可插拔的企业级中间件——跨引擎实例、跨请求、跨存储层级(GPU/CPU/本地磁盘/远程)复用任意位置(非仅前缀)的 KV,为 vLLM 生产部署提供了统一 cache 后端 [2510.09665]。Marconi 揭示了 hybrid(Attention+SSM/Mamba)模型下前缀缓存的根本性挑战——SSM 层的状态无法像 attention KV 那样按 token 前缀切分,需要 speculative insertion + FLOP-aware 驱逐策略,hybrid 模型 token hit rate 提升 4.5–34.4× [2411.19379]。生产 trace 分析进一步为 cache 设计提供实证依据——KVCache Cache in the Wild 刻画了云厂商大规模 KV cache 的真实访问模式(reuse distance、热度分布、workload 差异),提出 workload-aware 驱逐策略 [2506.02634]。KVDrive 将多层 KV 管理系统化——面向长上下文推理设计 GPU/CPU/SSD 多层协同的整体式 KV cache 管理,统一放置、预取与驱逐 [2605.18071]。Resident KV Claims 从契约角度重新定义 cache 复用——用"conformance contract"形式化"哪些 KV 承诺未来会被复用",在 active KV preservation 下管理常驻声明 [2605.24259]。跨模型维度上,ICaRus 发现不同模型间存在可复用的 identical cache,通过跨模型 KV 复用降低多模型 serving 的重复计算 [2603.13281];TokenDance 则在多 agent serving 中引入 collective KV cache sharing,将同一 workflow 内多 agent 的共享上下文 KV 集体复用 [2604.03143]

计算维度:CPU offloading 成为挖掘异构硬件潜力的新路线——FastDecode 将 attention 完全卸载到分布式远程 CPU,传输 QKV activation 而非 KVCache,实现 5.04× 吞吐 [2403.11421];NEO 选择更务实的本地 CPU asymmetric pipelining,以 greedy 回退保证永不劣于 baseline,T4 上达 7.5× [2411.01142]。硬件规格的代际提升未能自动转化为实际性能。MI300X 虽然理论 FP8 峰值 2615 TFLOPS 超过 H100 的 1979 TFLOPS,但实际利用率仅 45%(对比 NVIDIA 93%),根因是软件栈成熟度不足 [2510.27583]。Blackwell B200 通过 TMEM 和第五代 Tensor Core 实现了从训练专用到推理优化的架构转换,FP4 达到 96.3% 峰值利用率 [2512.02189] [2512.02189]。AMD CDNA4 (MI355X) 以 288 GB HBM3E 和 5 PFLOPS FP8 回应,但互联带宽落后 NVLink5 达 67% [amd-cdna4-whitepaper] [gpu-arch-comparison-2025]。BSP 编程模型在分布式执行中引入三类性能税——kernel launch、全局同步、数据局部性丧失——消耗掉潜在的并行加速 [2511.02168]。MI300X 的 chiplet 架构进一步引入 NUMA 效应:FlashAttention2 在 8-XCD 上 L2 命中率仅 ~1%,swizzled head-first mapping 可恢复至 90–97% 并加速 50% [2511.02132]。MoE 架构下,跨节点 all-to-all 通信几乎等于计算时间 [2412.19437]。Ops:Byte 比率从 A100 的 153 (FP16) 攀升到 MI355X/B200 的 ~315/281 (FP16) 和 629/563 (FP8),MXFP4 更达 1258 [gpu-arch-comparison-2025]——除大矩阵 GEMM 外几乎所有操作 memory-bound,FP4 的真正价值在于用更少 bit 传输相同语义信息。这一 memory-bound 认知的源头正是 FlashAttention 系列:FlashAttention (FA1) 首次将 attention 重构为 IO-aware 的 tiling + online softmax + recomputation,避免物化 $N^2$ 注意力矩阵,把 HBM 访问从 $O(N^2)$ 降到 $O(N)$,实现 2–4× 加速和 10–20× 内存节省 [2205.14135];FA2 通过重划分工作(减少非 matmul FLOPs、沿 seqlen 维并行、优化 warp 间通信)将 GPU 利用率从 25–40% 提升到 50–73% peak [2307.08691];FA3 引入 warp specialization + asynchrony(TMA/WGMMA 异步流水线)+ FP8 低精度,在 Hopper 上达到 75% peak(FP16)与近 1.2 PFLOPs/s(FP8)[2407.08608]——这条 FA1→FA2→FA3→FA4 谱系奠定了"attention kernel 随硬件代际重新设计"的方法论。

执行效率维度:FlashAttention-4 针对 Blackwell GPU 的非对称硬件 scaling(MMA 翻倍但 smem/exp 不变)重新设计 attention pipeline,利用 TMEM 和软件 exp 模拟在 B200 上达到 1613 TFLOPs/s (71% peak),超 cuDNN 1.3× [2603.05451]——标志注意力 kernel 从"通用高效实现"进入"硬件代际专用设计"阶段。即使存储和硬件不是瓶颈,kernel 间的执行边界本身也在消耗一个数量级的性能。8×H200 NVL 聚合内存带宽 ~38 TB/s,GLM-5.1 每 decode token 激活参数量 ~42 GB,纯带宽约束下理论吞吐可达 ~1000 tok/s,而实际系统仅输出几十 tok/s [tilert-speed-scaling-law]。根因不是算力不足,而是 GPU 反复执行 launch→load→compute→store→synchronize 循环,每次 kernel boundary 打断数据流、销毁局部性、强制重新同步 [tilert-speed-scaling-law]。这催生了三条并行的执行范式创新:TileRT 将整个 transformer AOT 编译为单个持久化 Engine Kernel [tile-ai-tilert],TokenSpeed 用 C++ FSM + Placement 编译器在 Blackwell 上超越 TRT-LLM [tokenspeed],Fleet 为 AMD MI350 多 die 设计了 hierarchical megakernel [ref:L2:2604.15379]。GPUOS 则走运行时路线——persistent kernel + lock-free ring buffer 将 per-op dispatch 从 3–7 μs 降至 <100 ns [2604.17861]。MPK 将 mega-kernel 思路产品化为通用编译器+运行时——把任意 tensor program 自动 mega-kernel 化为单个持久化 GPU kernel,通过编译期融合消除 kernel launch 与调度边界,把 TileRT/Fleet 的手工设计推广为可复用工具链 [2512.22219]。AVO 则从另一维度探索 kernel 优化的自动化——用 agentic variation operators 驱动的自主进化搜索来生成/优化 GPU kernel,把 kernel 设计从人工调优推向 LLM agent 自主演化 [2603.24517]。RTP-LLM 在模型加载环节展示了另一种执行效率优化——将加载范式从模型结构驱动重构为文件顺序 I/O,单进程读取 + 分布式 broadcast 消除冗余读取,TP=8 下实现 6.27× 加载加速(33s vs 204–207s),而 vLLM/SGLang 在同条件下出现负扩展性 [2605.29639] [2605.29639]

成本维度:FastServe 首次将 MLFQ 调度引入 LLM serving——利用 "input length 已知但 output length 未知" 的 semi information-agnostic 特性设计 skip-join 机制消除 HoL blocking,吞吐比 vLLM 提升 31.4× [2305.05920]。SLO 达标率与吞吐量之间的 Pareto 前沿尚未被有效推进。Colocation 方案在 P99 TBT 上表现糟糕,disaggregation 方案的 GPU 利用率低至 0.2% [2504.09285]。MoE disaggregated serving 面临 3 类网络流量(KV-cache reuse、collective 通信、P2D 传输)的竞争,传统 FIFO 调度无法区分流量优先级 [2603.17456]。ZeRO-Prefill 反转 MoE 数据流方向——从"路由激活到 expert"改为"流式传输 expert 权重到 GPU"——消除了 AllToAll 通信,MoE prefill 吞吐提升 1.35–1.59× [2605.02960]。配置空间爆炸使手动调优需要数天 GPU 时间 [2601.06288]。Justitia 引入 KV token-time(内存占用 × 时间)作为公平性度量,平均 JCT 减少 57.5% [ref:L2:2510.17015]。Harli 发现 decode 阶段 SM 利用率仅 ~40%,通过 inference+PEFT co-location 收割闲置算力实现 46% 额外推理吞吐——将"同 GPU 不同 workload 类型共存"作为新的利用率提升维度 [2511.11729]。FaaSMoE 走更极端的解耦路线——将 MoE expert 部署为 stateless FaaS 函数实现多租户 expert 共享和 scale-to-zero,尝试用 serverless 范式替代 dedicated GPU serving [2604.26881]。MARLIN 将优化目标从纯性能扩展到 sustainability——首次将碳排放、水消耗、电力成本与 TTFT 联合优化,在全球 geo-distributed DC 间做 game-theoretic routing [2605.13496]。Parallax 则打破了"管控集群"这一基本假设——在 volunteer WAN GPU pool 上用 decentralized DHT + two-phase DP 实现异构推理,3.6× 超越 HexGen [2509.26182]。RTP-LLM 的 cache 感知流量调度从成本角度展示了生产级优化的极致效果——75% prefill 机器缩减(80 → 20 台)直接转化为运营成本节约,35–37% TTFT P95 降低无吞吐牺牲 [2605.29639]。MoE 的运行时负载均衡则从 dispatch 粒度收割成本——SGLang 的 Waterfill 把 dense shared expert 当作可分派 slot 按各 rank 负载"填谷"(V3/R1 +1.48%~+4.66%),LPLB 用每层 min–max LP 在 redundant 副本间实时重分流量(红/绿双向,最高 +7.34%),两者均不改 logical top-k 保持语义 [blog-waterfill-lplb]。PD 解离下的 MoE decode 路由进一步引入 expert 局部性维度——ELDR 发现 decode 实例选择应感知 expert 放置局部性,为 PD-disaggregated MoE serving 设计 expert-locality-aware 的 decode routing,减少跨节点 expert 访问 [2607.00466]

RDMA 调度器解耦维度:一条新的技术轴正在成形——将调度与协调功能从 CPU control plane 迁移到 RDMA fabric 和 SmartNIC 硬件。CPU-Slowdowns 的系统性测量揭示了 multi-GPU 推理中 CPU 的隐藏主导瓶颈:tokenization 占 TTFT 高达 50%,shared-memory broadcast dequeue 延迟膨胀 19×(12 ms → 228 ms),NCCL barrier 同步将单核延迟放大为全局 GPU 停顿 [2603.22774]。生产集群日志(4.65M 条 salloc 记录)显示 H100 节点 P25 的 CPU-to-GPU ratio 仅为 0.25——1 核服务 4 GPU [2603.22774]。Blink 率先证明 CPU 可被完全移除:SmartNIC (BlueField-3 DPU) 承担请求管理和 RDMA 传输,GPU-resident persistent scheduler(单个 256 线程 CUDA kernel 无限运行)接管调度与执行,P99 TTFT 降低最高 8.47×,在 CPU 干扰下保持 0.92–1.14× 稳定而 baselines 退化 1.54–18.84× [2604.07609]。这一架构的数据通路完全依赖 RDMA:DPU 通过 one-sided RDMA write 将 tokenized prompts 写入 GPU-resident ring buffer,GPU scheduler 通过 atomic CAS 声明 slot 所有权 [2604.07609]。OnePiece 将类似的 RDMA ring buffer 模式扩展到 AIGC 多阶段推理流水线——one-sided RDMA 传输中间结果,double-ring buffer 解决无 CPU 参与的 RDMA 死锁,Node Manager 动态调度 GPU 实例 [2601.20655]。这些系统的 RDMA 数据面根植于更早的基础研究:FaRM 建立了 RDMA-write circular-buffer messaging 和 lock-free one-sided reads 的范式,实现 10× 吞吐和 100× 低延迟(vs TCP/IP),证明 RDMA 可作为分布式系统的一等通信原语 [farm-nsdi14]。KRCore 解决了 RDMA 的弹性计算瓶颈——通过 DCT 虚拟化将连接建立从 15.7ms 降至 5.4μs(2,900×),使 RDMA 适用于 serverless/elastic 场景中的海量短连接 [2201.11578]。RackSched 则在网络层面探索了微秒级调度的硬件化——将 ToR 可编程交换机用作 rack 级调度器,power-of-k-choices 做服务器间负载均衡,实现近线性吞吐扩展并保持单服务器水平的尾延迟 [racksched-osdi20]

这六个维度的矛盾在 2022–2026 年同时激化,催生了本主题 84 篇跨 framework/hardware/cluster/kernel 四大类别的系统性工作——从 vLLM/PagedAttention 奠基 [2309.06180] 到 Blink 的 CPU-free 架构变革 [2604.07609],再到 RTP-LLM 以生产系统实证验证 cache-centric + PD 解离 + 模块化推测解码的全栈集成方案 [2605.29639],覆盖了存储到执行、单机到全球、CPU-centric 到 RDMA-centric 的完整技术栈。

Agentic 基础设施维度:Agentic workflow 将单个用户查询拆成 19–21 次 LLM 调用 [2505.05286],但现有独立请求调度器既不懂阶段依赖也不懂异构 GPU 分配。Pie (SOSP 2025) 从根本上质疑 monolithic 生成循环——将 prefill-decode 分解为 embed/forward/sample 可编程 handler,通过 Wasm inferlets 将控制权交给应用层,agentic 吞吐 1.3–3.4× [2510.24051]。FlowMesh 在更高层次用 DAG operator identity + utility-based scheduling 实现跨租户去重和异构 GPU 调度,成本降低 1.8–3.8× [2510.26913]。GORGO 将视野扩展到跨地理区域——用 additive cost model 联合优化 network latency + prefix overlap + queue depth,跨区域 TTFT 2.5× [2602.11688]。高并发 agentic batch inference 中 KV-Cache thrashing 使中间阶段 49% 的延迟来自 KV 重计算,Concur 用 AIMD 准入控制将 KV 命中率从 35% 稳定在 73–96% [ref:L2:2601.22705]。ThunderAgent 引入"agentic program"作为一等调度单元,实现 1.48–3.58× serving 吞吐提升 [ref:L2:2602.13692]。Helium 借鉴数据库查询优化器思路,用 Templated Radix Tree 做 plan-level prefix reuse 优化 [ref:L2:2603.16104]。Halo 将 batch agentic workflow 整合为 DAG 统一调度,比 naive vLLM 加速 400× [ref:L2:2509.02121]。这一波 agentic serving 的前缀复用范式源自 SGLang——它用 RadixAttention(radix tree 组织 KV cache)实现自动跨请求前缀复用,配合 compressed FSM 约束解码和协同调度,为结构化 LLM 程序提供了 5× 吞吐提升,是几乎所有后续 prefix-caching agentic 系统的基础 [2312.07104]。工具型 agent 的调度则需要 orchestrator 与 engine 协同——Sutradhara 提出 orchestrator-engine co-design,让上层 agent 编排器与底层推理引擎交换阶段/依赖信息,为 tool-based agentic 推理做智能调度 [2601.12967]。GLM Graph-CoT 则针对图链式推理这一新 agentic 模式——多 agent 框架配合高效 serving,将 graph chain-of-thought reasoning 规模化 [2511.01633]

推测解码与执行加速维度:推测解码正从算法研究走向系统级调度。SSD(Speculative Speculative Decoding)对推测过程本身再做推测,进一步压缩草稿开销 [2603.03251]。SPECTRE 引入 hybrid ordinary-parallel speculative serving,在资源受限下混合普通解码与并行推测以提升资源效率 [2605.08151]。DSpark 把"起草更快更准"与"验证更省"合成一个无损系统——半自回归草稿(并行主干 + 轻量序列头注入块内依赖)+ 置信度调度验证(按实测 SPS(B) 吞吐曲线动态选每请求验证长度),在 DeepSeek-V4 上较 MTP-1 每用户生成加速 57–85%,把服务吞吐-交互 Pareto 前沿外推 [dspark]。CPU-GPU 协同执行的另一路线由 APEX 代表——在资源受限设备上做异步并行 CPU-GPU 执行,重叠 CPU 与 GPU 阶段以提升在线推理效率 [2506.03296]。FlowKV 则从 PD 传输角度补齐——为 disaggregated inference 设计低延迟 KV cache 传输框架,降低 prefill→decode 的 KV 搬运延迟 [2504.03775]。NetKV 将网络感知引入 decode 实例选择——为 disaggregated LLM inference 做 network-aware 的 decode instance selection,在 KV 传输代价与算力负载间联合优化 [2606.03910]

§2 覆盖的 category 分布 #

Category论文数代表性工作
framework61vLLM [2309.06180]、Mooncake [2407.00079]、SGLang [2312.07104]、DualPath [2602.21548]、PrfaaS [2604.15039]、Blink [2604.07609]、FastServe [2305.05920]、Pie [2510.24051]、FlowMesh [2510.26913]、TileRT [tile-ai-tilert]、ZeRO-Prefill [2605.02960]、RTP-LLM [2605.29639]、CacheGen [2310.07240]、LMCache [2510.09665]、Marconi [2411.19379]、FlowKV [2504.03775]、KVDrive [2605.18071]、ICaRus [2603.13281]、TokenDance [2604.03143]、DSpark [dspark]、SSD [2603.03251]、SPECTRE [2605.08151]、ELDR [2607.00466]、Sutradhara [2601.12967]
hardware4MI300X 性能分析 [2510.27583]、Blackwell 微架构 [2512.02189]、GPU 架构对比 [gpu-arch-comparison-2025]、CDNA4 白皮书 [amd-cdna4-whitepaper]
cluster9Multi-GPU 性能税 [2511.02168]、fabric-lib [2510.27656]、MARLIN [2605.13496]、KRCore [2201.11578]、RackSched [racksched-osdi20]、OnePiece [2601.20655]、FaRM [farm-nsdi14]、CPU-Slowdowns [2603.22774]、NetKV [2606.03910]
kernel10ConCCL DMA Offload [2412.14335]、GPUOS [2604.17861]、NUMA-Attn [2511.02132]、HipKittens [2511.08083]、FlashAttention-4 [2603.05451]、FlashAttention [2205.14135]、FlashAttention-2 [2307.08691]、FlashAttention-3 [2407.08608]、MPK [2512.22219]、AVO [2603.24517]

本轮新增 RTP-LLM(framework 类别)使 framework 论文数达到 42 篇。RTP-LLM 作为阿里巴巴服务 100M+ 用户的生产级推理引擎,横跨 PD 解离、KV cache 管理、推测解码、多模态服务和模型加载五个技术维度 [2605.29639]。其 cache 感知流量调度 + 4 层 KV cache 层级 + 文件序加载的全栈集成设计是 cache-centric serving 范式的完整生产验证 [2605.29639]

本次增量新增 25 篇(含 1 篇工程 blog)后总数达 84。其中 kernel 类别显著扩张——补齐了 FlashAttention 完整谱系(FA1/FA2/FA3)作为 FA4 的前置演进,并新增 MPK(通用 mega-kernel 编译器)和 AVO(agentic kernel 进化搜索)。framework 类别新增大量 KV cache 管理(CacheGen/LMCache/Marconi/KVDrive/ICaRus/TokenDance/Resident KV Claims/KV survey)、PD 解离与 MoE 路由(FlowKV/ELDR/NetKV)、推测解码(SSD/SPECTRE/DSpark)与 agentic 编排(SGLang/Sutradhara/GLM Graph-CoT)工作,标志 KV-cache-centric 与推测解码两条技术线的系统化成熟。

§3 时间线 #

时间里程碑关键贡献
2014FaRMRDMA-write circular-buffer messaging + lock-free one-sided reads,10× 吞吐 / 100× 低延迟 vs TCP/IP,奠定 RDMA 分布式系统通信范式 [farm-nsdi14]
2020RackSchedToR 可编程交换机用作 rack 级 μs 调度器,power-of-k-choices + tree-based parallel min,近线性吞吐扩展(OSDI 2020)[racksched-osdi20]
2022-01KRCoreDCT 虚拟化将 RDMA 连接建立从 15.7ms 降至 5.4μs(2,900×),O(1) 内存,支持 elastic/serverless 场景 [2201.11578]
2022-05FlashAttentionIO-aware tiling + online softmax + recomputation,HBM 访问 $O(N^2)\to O(N)$,2–4× 加速 / 10–20× 内存节省,奠定 memory-bound attention 范式 [2205.14135]
2023-05FastServeSkip-join MLFQ + proactive KV swap 消除 HoL blocking,吞吐比 vLLM 提升 31.4× [2305.05920]
2023-07FlashAttention-2减少非 matmul FLOPs + 沿 seqlen 并行 + 优化 warp 通信,GPU 利用率 25–40%→50–73% peak [2307.08691]
2023-09vLLM/PagedAttentionOS 分页思想引入 KV cache 管理,block table + CoW 消除碎片与冗余复制,吞吐提升 2–4× [2309.06180]
2023-10CacheGenKV cache 算术编码 + 分层量化,3.5–4.3× 压缩 + 带宽自适应流式加载,TTFT 降低 3.2–3.7× [2310.07240]
2023-12SGLangRadixAttention(radix tree KV 复用)+ compressed FSM 约束解码 + 协同调度,结构化 LLM 程序 5× 吞吐,奠定 prefix-caching 范式 [2312.07104]
2023-12FlexRLHFDisaggregated 策略将 RLHF 推理与训练物理分离,65B 模型 11× 加速 [2312.11819]
2024-03FastDecodeS-Part/R-Part 分解将 KVCache 留在分布式 CPU,传输 QKV activation 替代 KVCache,5.04× over vLLM [2403.11421]
2024-05DeepSeek-V2MLA KVCache 压缩 93.3% + DeepSeekMoE 160 experts,21B 激活达 70B dense 水平,推理 5.76× [2405.04434]
2024-07MooncakeKVCache-centric disaggregated 架构 + 分布式 DRAM Pool + CPP 长文本 pipeline,比 vLLM 多处理 75% 请求 [2407.00079]
2024-07FlashAttention-3Warp specialization + asynchrony(TMA/WGMMA)+ FP8,Hopper 上 75% peak(FP16)/ 近 1.2 PFLOPs/s(FP8)[2407.08608]
2024-09HybridFlow/veRLHybrid single/multi-controller + 3D-HybridEngine 实现零冗余 resharding,1.53–20.57× RLHF 加速 [2409.19256]
2024-11NEOAsymmetric GPU-CPU pipelining 卸载 decode attention 到本机 CPU,greedy 回退保证不劣于 baseline,T4 7.5× [2411.01142]
2024-11MarconiHybrid(Attention+SSM)前缀缓存:speculative insertion + FLOP-aware 驱逐,hybrid 模型 token hit rate 4.5–34.4× [2411.19379]
2024-12DeepSeek-V3DualPipe 双向 pipeline 实现 MoE all-to-all 通信完全隐藏,671B 模型训练仅 $5.576M [2412.19437]
2024-12ConCCL首次在不改硬件前提下通过 DMA offload 将 C3 加速比从 21% 提升至 72% [2412.14335]
2024-12KV Cache SurveyKV cache 管理综述:token-level / model-level / system-level 三层分类法,统一问题框架 [2412.19442]
2025-01OnePieceOne-sided RDMA + double-ring buffer 实现 AIGC 多阶段微服务间零 CPU 通信,Node Manager 动态 GPU 调度 [2601.20655]
2025-04DynaServe微请求抽象统一 colocation 和 disaggregation,容量提升 1.15–3.07× [2504.09285]
2025-04JITServe首个面向多类型 SLO 的调度系统,goodput 提升 1.4–6.3×,常数竞争比 1/8.55 [2504.20068]
2025-04FlowKVDisaggregated inference 低延迟 KV cache 传输框架,降低 prefill→decode KV 搬运延迟 [2504.03775]
2025-06KVCache in the Wild云厂商大规模 KV cache 访问模式刻画(reuse distance/热度/workload 差异)+ workload-aware 驱逐 [2506.02634]
2025-06APEX资源受限设备上异步并行 CPU-GPU 执行,重叠 CPU/GPU 阶段提升在线推理效率 [2506.03296]
2025-05HexGen-Flow两层调度器处理 agentic workflow + 异构 GPU,P95 延迟降低 1.42–1.56× [2505.05286]
2025-07KVFlowWorkflow-aware 的 STE 驱逐策略替代 LRU,agentic 场景加速 1.83–2.19× [2507.07400]
2025-09Halo首个将数据库 batch query optimization 应用于 agentic workflow,比 naive vLLM 加速 >400× [ref:L2:2509.02121]
2025-09Parallax首个面向 volunteer WAN GPU pool 的去中心化推理系统,two-phase DP 分解 + cross-replica stitching,3.6× over HexGen [2509.26182]
2025-10PieWasm inferlets 分解 monolithic 生成循环为可编程 handler,agentic 吞吐 1.3–3.4×(SOSP 2025)[2510.24051]
2025-10FlowMeshDAG operator + 双层 hash identity ($H_{task}$/$H_{exec}$) + utility-based 异构调度,成本降低 1.8–3.8× [2510.26913]
2025-10MI300X 性能画像首个系统性量化:实际利用率 45%,软件效率 80–85%,频率降低 ~42% [2510.27583]
2025-10JustitiaVirtual-Time Fair Queuing + KV token-time 公平性度量,JCT 减少 57.5% [ref:L2:2510.17015]
2025-10fabric-lib可移植 RDMA P2P 库统一 ConnectX-7 和 AWS EFA,MoE decode 超越 DeepEP [2510.27656]
2025-10LMCache企业级可插拔 KV cache 层:跨引擎/请求/存储层级复用任意位置 KV,vLLM 生产统一 cache 后端 [2510.09665]
2025-11GLM Graph-CoT多 agent 框架 + 高效 serving 规模化 graph chain-of-thought reasoning [2511.01633]
2025-11Three TaxesBSP 性能税分析框架 + Iris compute-comm fusion,Flash Decode 加速 10–20% [2511.02168]
2025-11NUMA-AttnSwizzled head-first mapping 恢复 MI300X L2 命中率至 90–97%,50% attention 加速 [2511.02132]
2025-11HipKittens首个 AMD 系统化 tile-based C++ DSL,BF16 GEMM 达 1610 TFLOPS,匹配手写汇编 [2511.08083]
2025-11Harli首个 inference+PEFT co-location 系统,GreenContext SM 分区 + CUDA VMM 统一分配器,46% 额外推理吞吐 [2511.11729]
2025-12Blackwell 微基准TMEM 延迟降低 58%,tcgen05 单指令延迟 11 cycles vs Hopper 32–128 cycles,FP32 accumulator 使 TC 吞吐减半 [2512.02189]
2025-12MoE 负载均衡ILP + 多项式时间启发式联合优化负载均衡与数据搬运开销 [moe-lb-interai25]
2025-12MPK通用 mega-kernel 编译器 + 运行时,自动将任意 tensor program mega-kernel 化为单持久化 kernel,消除 launch/调度边界 [2512.22219]
2026-01AIConfigurator无 GPU profiling 配置搜索,比实际 benchmarking 快 170,000–427,000× [2601.06288]
2026-01ConcurAIMD 准入控制将 KV 命中率从 35% 稳定在 73–96%,Qwen3-32B 吞吐提升 4.09× [ref:L2:2601.22705]
2026-01SutradharaOrchestrator-engine co-design,上层编排器与推理引擎交换阶段/依赖信息,tool-based agentic 智能调度 [2601.12967]
2026-02GORGOAdditive cost model 联合优化 network latency + prefix overlap + queue depth,跨区域 TTFT 2.5× [2602.11688]
2026-02DualPath聚合 DE 空闲 SNIC 带宽实现 1.87× 离线吞吐提升,1152 GPU 近线性扩展 [2602.21548]
2026-02ThunderAgent"Agentic program"一等调度单元,RL rollout 吞吐提升 1.79–3.92× [ref:L2:2602.13692]
2026-03PPDAppend-prefill 仅 2% TPOT 劣化,Turn 2+ TTFT 降低 48–73% [2603.13358]
2026-03FlashAttention-4TMEM pipeline + 软件 exp 模拟 + 2-CTA MMA,B200 BF16 达 1613 TFLOPs/s (71% peak),1.3× over cuDNN [2603.05451]
2026-03MFSRMLQ 三阶段流量调度解决 disaggregated MoE 网络竞争,SLO 达标率提升 1.2–2.4× [2603.17456]
2026-03HeliumTemplated Radix Tree + DB-style plan 优化,超越 KVFlow 1.56× [ref:L2:2603.16104]
2026-03CPU-Slowdowns首次系统性量化 multi-GPU 推理中 CPU 隐藏瓶颈:tokenization 占 TTFT 50%,broadcast 19× 膨胀,barrier 同步放大 [2603.22774]
2026-03SSDSpeculative Speculative Decoding:对推测过程本身再做推测,进一步压缩草稿开销 [2603.03251]
2026-03ICaRus跨模型 identical cache 复用,降低多模型 serving 重复计算 [2603.13281]
2026-03AVOAgentic variation operators 驱动的自主进化搜索生成/优化 GPU kernel [2603.24517]
2026-04PrfaaS跨 DC PD 分离,hybrid attention 使 KV 吞吐降至 3 Gbps,+54% 吞吐 [2604.15039]
2026-04TensorHubROS 无所有权存储 + pipeline replication,RL 权重传输 stall 减少 6.7× [2604.09107]
2026-04BlinkCPU-free LLM inference:SmartNIC frontend + GPU-resident persistent scheduler,P99 TTFT 降低最高 8.47×,CPU 干扰下 baselines 退化 18.84× 而 Blink 保持稳定 [2604.07609]
2026-04ScepsyAggregate LLM Pipeline 抽象 + 分数 GPU 调度,吞吐提升 2.4× [2604.15186]
2026-04FleetChiplet-task 抽象 + hierarchical megakernel,MI350 上 decode 加速 1.3–1.56× [ref:L2:2604.15379]
2026-04ForkKVDualRadixTree + ResidualAttention 分离 base/residual KV cache,multi-LoRA agent 吞吐 1.25–3.04× [2604.06370]
2026-04FaaSMoEMoE experts-as-FaaS-functions 架构解耦 + 多租户 expert 共享,<1/3 资源消耗 [2604.26881]
2026-04TokenDanceMulti-agent serving 中 collective KV cache sharing,同 workflow 多 agent 共享上下文 KV 集体复用 [2604.03143]
2026-04GPUOSPersistent kernel + dynamic operator injection,micro-ops 加速 15.3×,能耗降低 20–22% [2604.17861]
2026-05TileRTAOT 编译全模型为单个 Persistent Engine Kernel,warp/block/GPU 三级特化,GLM-5.1 生产上线 [tile-ai-tilert]
2026-05TokenSpeedC++ FSM + Placement 编译器 + Blackwell MLA 内核,超越 TRT-LLM 9–11% [tokenspeed]
2026-05ZeRO-PrefillAsyncEP 反转 MoE 数据流方向,消除 AllToAll,MoE prefill 吞吐提升 1.35–1.59× [2605.02960]
2026-05KVServeService-aware KV 压缩 + controller,最高 10× KV 数据削减 [kvserve]
2026-05MARLIN首个 sustainability-aware geo-distributed LLM 推理调度器,multi-agent game-theoretic RL,PHV 46.4% 超越次优 [2605.13496]
2026-05RTP-LLM阿里生产级推理引擎(100M+ 用户),集成 PD 解离 + 4 层 KV cache 层级 + cache 感知调度 + 模块化推测解码 + 文件序加载 + EPD 多模态解离,480B MoE 上 4.72–5.33× TTFT,6.27× 加载加速 [2605.29639]
2026-05SPECTREHybrid ordinary-parallel speculative serving,资源受限下混合普通解码与并行推测提升资源效率 [2605.08151]
2026-05KVDrive面向长上下文的整体式多层 KV cache 管理(GPU/CPU/SSD 协同放置/预取/驱逐)[2605.18071]
2026-05Resident KV ClaimsConformance contract 形式化"承诺未来复用的 KV",active KV preservation 下管理常驻声明 [2605.24259]
2026-06NetKVNetwork-aware decode instance selection,KV 传输代价与算力负载联合优化 [2606.03910]
2026-07ELDRExpert-locality-aware decode routing,PD-disaggregated MoE serving 减少跨节点 expert 访问 [2607.00466]
2026-07DSpark半自回归草稿 + 置信度调度验证的无损投机解码,DeepSeek-V4 上每用户生成加速 57–85%,外推吞吐-交互 Pareto 前沿 [dspark]
2026-07Waterfill + LPLBSGLang dispatch-time MoE 负载均衡:Waterfill 填谷 shared expert(+1.48%~+4.66%)+ LPLB min–max LP 副本重分流(最高 +7.34%),不改 logical top-k [blog-waterfill-lplb]

§4 Evolution timeline (技术谱系) #

flowchart TB subgraph Storage["存储与 KV-Cache 演进"] direction TB vLLM_PA["vLLM/PagedAttention: block table + CoW
(2023) framework"] --> MC["Mooncake: 分布式 DRAM Pool
(2024) framework"] vLLM_PA --> vLLM_PD["vLLM PD split"] MC --> DP["DualPath: 聚合空闲 SNIC
framework+cluster"] MC --> PrfaaS["PrfaaS: 跨 DC KV 传输
framework"] MC --> RTPLLM["RTP-LLM: 4层KV层级
+统一hash map framework"] DSv2_MLA["DeepSeek-V2 MLA: KV 压缩 93.3%
(2024) framework"] --> PrfaaS DP --> KVF["KVFlow: workflow-aware 驱逐
framework"] PrfaaS --> KVS["KVServe: service-aware
KV 压缩 framework"] KVF --> Helium["Helium: TRT + plan 优化
framework"] KVF --> Concur["Concur: AIMD 准入控制
framework"] NCCL["NCCL Broadcast"] --> TH["TensorHub: ROS 无所有权引用
framework+cluster"] UCX["UCX P2P"] --> TH vLLM_PD --> PPD["PPD: append-prefill 路由
framework"] vLLM_PA --> ForkKV_n["ForkKV: DualRadixTree
base/residual KV 分离 framework"] vLLM_PA --> RTPLLM RTPLLM -.->|"cache-centric 范式
共同验证"| MC end subgraph Compute["计算与执行演进"] direction TB Hopper["Hopper H100/H200
hardware"] --> BWell["Blackwell B200: TMEM+tcgen05
hardware"] BWell --> FA4["FlashAttention-4: TMEM pipeline
+ 软件 exp 模拟 kernel"] CDNA3["MI300X CDNA3
hardware"] --> Bench["MI300X 45% 利用率诊断
hardware"] Bench --> CDNA4["MI355X CDNA4: MXFP4 10PF
hardware"] CDNA3 --> NUMA["NUMA-Attn: swizzled mapping
kernel"] CDNA4 --> HK["HipKittens: tile DSL
kernel"] BSP["BSP 编程模型"] --> Taxes["Three Taxes 分析
cluster+kernel"] Taxes --> Iris["Iris compute-comm fusion
kernel"] Iris --> TileRT["TileRT: 全模型 AOT
Engine Kernel framework"] BWell --> TS["TokenSpeed: FSM+Placement
framework"] CDNA4 --> Fleet["Fleet: chiplet megakernel
kernel"] RCCL["RCCL/NCCL Collectives"] --> FL["fabric-lib: P2P RDMA
cluster"] FL --> Iris DSv3_Comm["DeepSeek-V3 warp spec
framework+kernel"] --> Taxes DSv3_Comm --> TileRT Launch["kernel launch 3-7μs"] --> GPUOS["GPUOS: persistent+ring buffer
kernel"] end subgraph RDMASchedDisagg["RDMA 调度器解耦演进"] direction TB FaRM_n["FaRM: RDMA-write circular buffer
+ lock-free reads (2014) cluster"] FaRM_n --> KRCore_n["KRCore: DCT 虚拟化
5.4μs 连接建立 (2022) cluster"] FaRM_n --> OP["OnePiece: RDMA double-ring
buffer AIGC (2025) cluster"] KRCore_n --> OP CPUSlow["CPU-Slowdowns: CPU 隐藏瓶颈
50% TTFT · 19× IPC 膨胀 (2026) cluster"] CPUSlow --> Blink_n["Blink: CPU-free inference
DPU + persistent GPU sched (2026) framework"] OP --> Blink_n RS["RackSched: ToR μs scheduler
power-of-k-choices (2020) cluster"] RS -.->|"in-network sched concept"| Blink_n FaRM_n -.->|"RDMA messaging primitive"| FL KRCore_n -.->|"elastic connection"| FL end subgraph CPUOffload["CPU Offloading 演进"] direction TB FD["FastDecode: S/R-Part 分解
远程 CPU 集群 framework"] --> NEO_n["NEO: asymmetric 本地 CPU
pipelining framework"] MC -.->|"KV 存储于 CPU DRAM"| NEO_n CPUSlow -.->|"motivation: CPU is bottleneck"| NEO_n end subgraph Cost["成本与调度演进"] direction TB FastServe_n["FastServe: skip-join MLFQ
(2023) framework"] --> DS["DynaServe: 微请求抽象
framework"] vLLM_FCFS["vLLM FCFS"] --> FastServe_n DS --> JIT["JITServe: 多类型 SLO
framework"] DS --> HF["HexGen-Flow: 异构 GPU
framework"] JIT --> Scp["Scepsy: 分数 GPU pipeline
framework"] HF --> Scp RR["Round-Robin EP"] --> MoELB["MoE 负载均衡 ILP
framework"] DSv3_PP["DeepSeek-V3 DualPipe
framework"] --> DS Bench2["GPU Arch Comparison
hardware"] --> AIC["AIConfigurator
framework"] DS --> MFS_n["MFS: RMLQ 三阶段流量
framework"] MoELB --> ZP["ZeRO-Prefill: AsyncEP
framework"] WFQ["WFQ 网络调度"] --> Just["Justitia: KV token-time
framework"] HF --> Parallax_n["Parallax: decentralized WAN
framework"] DS --> Harli_n["Harli: inference+PEFT co-location
framework"] ZP --> FaaSMoE_n["FaaSMoE: expert-as-FaaS
framework"] RTPLLM_cost["RTP-LLM: cache-aware
75% prefill节省 framework"] --> DS end subgraph GeoDist["Geo-Distributed 演进"] direction TB GORGO_n["GORGO: additive cost model
跨区域 KV routing framework"] --> MARLIN_n["MARLIN: sustainability-aware
inter-DC routing cluster"] PrfaaS_gd["PrfaaS 跨 DC prefill"] --> MARLIN_n PrfaaS_gd --> GORGO_n MFS_gd["MFS 网络调度"] -.-> MARLIN_n end subgraph Agentic["Agentic 工作负载演进"] direction TB K8s["K8s + vLLM"] --> TA["ThunderAgent: program-aware
framework"] K8s --> Halo_n["Halo: batch DAG 优化
framework"] K8s --> Pie_n["Pie: Wasm inferlets 可编程 serving
(SOSP 2025) framework"] K8s --> FM_n["FlowMesh: DAG fabric + 异构调度
framework"] KVF --> Helium TA --> Concur Pie_n -.->|"handler-level vs
workflow-level"| FM_n end subgraph SpecDecode["推测解码演进"] direction TB DSv3_MTP["DeepSeek-V3 MTP
framework"] --> RTPLLM_SD["RTP-LLM: 模块化推测解码
Naive/MTP/Eagle/PromptLookup framework"] Medusa["Medusa heads
algorithm"] --> RTPLLM_SD Eagle["Eagle speculative
algorithm"] --> RTPLLM_SD end DP -.->|"RDMA 带宽聚合
共同洞察"| TH FL -.->|"传输引擎"| TH DP -.->|"存储 I/O 约束
驱动调度"| DS BWell -.->|"FP4/FP6 新精度
改变 ops:byte"| AIC MoELB -.->|"expert 放置优化"| Scp Taxes -.->|"inter-kernel idle 诊断"| TileRT GPUOS -.->|"JIT 动态注入 vs
AOT 静态编译"| TileRT Fleet -.->|"chiplet megakernel vs
NVLink persistent"| TileRT NUMA -.->|"chiplet locality"| Fleet HK -.->|"AMD kernel 基座"| Fleet Parallax_n -.->|"decentralized vs
centralized PD"| DP MARLIN_n -.->|"sustainability vs
throughput-max"| JIT FaaSMoE_n -.->|"serverless vs
dedicated GPU"| ZP Harli_n -.->|"SM 收割 vs
SNIC 收割"| DP FA4 -.->|"Blackwell kernel vs
CDNA4 kernel"| Fleet DSv2_MLA -.->|"MLA 驱动
KV 压缩"| KVS ForkKV_n -.->|"multi-LoRA CoW vs
单模型 CoW"| vLLM_PA GORGO_n -.->|"routing cost model
vs game-theoretic"| MARLIN_n NEO_n -.->|"local CPU vs
remote CPU"| FD FastServe_n -.->|"MLFQ scheduling
foundation"| Just Blink_n -.->|"CPU-free vs
CPU-offload"| NEO_n Blink_n -.->|"persistent GPU sched
same primitive"| GPUOS CPUSlow -.->|"broadcast 竞争
motivates DPU"| Blink_n RTPLLM -.->|"cache-aware调度 vs
CPU-free调度"| Blink_n RTPLLM_SD -.->|"C++实现消除
Python开销"| GPUOS

关键谱系节点(含本轮新增):

§5 技术线交错 #

5.1 存储选择如何约束计算架构 #

KV-Cache 的存取方式直接决定了 GPU 的有效利用率。DualPath 的 CNIC-centric 数据通路设计本质上是对"当前 GPU 不支持 PCIe QoS"这一硬件缺陷的软件 workaround [2602.21548]。PrfaaS 的层间 prefill pipelining 要求 execution engine 能在每层计算完成后立即发起 KV 传输 [2604.15039]。KVServe 的 service-aware 压缩在 PD 通路上引入了新的 trade-off:10× KV 压缩使低带宽链路可行,但需要 decode 端的解压开销 [kvserve]。PPD 则揭示了一个更精细的约束——append-prefill 与 full prefill 的干扰模式完全不同,前者仅 2% TPOT 劣化,使 Turn 2+ 可安全 co-locate 于 decode 节点 [2603.13358]。RTP-LLM 的四层 cache 层级引入了另一重约束——调度决策必须同时考虑 cache 所在层级的访问代价:GPU 驻留 cache 零传输、本地 CPU 需 PCIe 拷贝、远程 CPU 需 RDMA 传输、3FS 需跨网络 fetch,每层的延迟差异直接体现在评分函数的 $\alpha > \beta$ 权重差中 [2605.29639] [2605.29639]

5.2 硬件能力如何开启新框架设计 #

MI300X 的 192 GB HBM 允许更大模型单卡部署,减少 tensor parallelism 需求 [2510.27583]。MI355X 将显存扩展到 288 GB、FP8 算力到 5 PFLOPS,MXFP4 达 10 PFLOPS [amd-cdna4-whitepaper]。Blackwell 的 FP4 达到 9000 TFLOPS(实测 96.3% 峰值利用率 [2512.02189]),但 FP32 accumulator 使 Tensor Core 吞吐减半。NVLink5 (1.8 TB/s) + NVL72 (72-GPU 域 130 TB/s) 与 Infinity Fabric (1075 GB/s, 8-GPU 域) 的 67% 带宽差距 [amd-cdna4-whitepaper] 迫使 fabric-lib 设计跨 vendor 通信抽象 [2510.27656]

5.3 模型架构如何强制部署重新思考 #

DeepSeek-V3 的 256 routed experts + 64-way EP 使跨节点 all-to-all 通信时间几乎等于计算时间 [2412.19437]。DualPipe 通过双向调度和 warp specialization 将 20/132 SM(15%)专门分配给通信才能实现完全隐藏 [2412.19437]。ZeRO-Prefill 走更激进的路线——完全消除 AllToAll,用背景 AllGather 流式传输 expert 权重 [2605.02960]。RTP-LLM 在 480B MoE 模型(Qwen3-Coder)上的 PD 解离部署验证了大规模 MoE 的生产可行性——5 节点 × 8 GPU,使用 DeepEP 做 All2All 通信,4.72–5.33× TTFT 加速且吞吐不牺牲 [2605.29639]

5.4 持续执行范式如何跨越 framework-kernel-hardware 边界 #

TileRT 的 Persistent Engine Kernel 是系统设计和 GPU 编程模型深度融合的产物——AOT 编译器将整个 transformer 模型静态展开为单个 kernel [tile-ai-tilert]。GPUOS 走最保守的路线——保留 PyTorch eager semantics,但在 persistent kernel 中用 lock-free ring buffer 接收动态注入的 operator [2604.17861]。Blink 的 GPU-resident persistent scheduler 是 GPUOS 思路在 serving 层面的极致应用——不是持久化一个计算 kernel,而是持久化整个调度循环(scan ring buffer → claim slot → select graph → launch → poll → publish),永久占据 1 个 thread block 的 256 线程 [2604.07609]。TileRT/GPUOS/Blink 三者共同验证了一个原则:GPU-resident 执行的真正价值不在于减少单次 launch latency(2μs vs 11μs),而在于将整个 host-device 交互模式从"per-step round-trip"转变为"persistent polling"。RTP-LLM 的模块化 C++ 推测解码框架 [2605.29639] 虽未走 persistent kernel 路线,但通过消除 Python-to-C++ 调用开销(对 vLLM 1.12× 吞吐提升 [2605.29639])体现了同一精神——减少执行边界处的软件栈开销。

5.5 调度器与通信库的纵向协同 #

fabric-lib 证明 P2P 通信是 collective 的必要补充——collectives 的四大约束(fixed membership、synchronized initialization、operation ordering、shape uniformity)在 disaggregated inference/MoE 中已成瓶颈 [2510.27656]。TensorHub 的 ROS 利用 RDMA 引擎实现零 trainer stall [2604.09107]。FaRM 更早地揭示了 RDMA 编程模型的核心 trade-off——one-sided operations 绕过远端 CPU 实现最大吞吐但要求本地版本验证,messaging(RDMA-write ring buffer)保留语义灵活性但消耗远端 CPU polling [farm-nsdi14]。KRCore 进一步发现当远端 CPU 仅需提供连接元数据(12B/node)时,one-sided RDMA READ 可完全绕过 meta server CPU 实现 11.8× 吞吐 [2201.11578]——这一 insight 直接适用于 LLM serving 中 scheduler metadata 的分发。RTP-LLM 的 4 层 cache 层级中第三层(远程 CPU RDMA)[2605.29639] 正是将 RDMA P2P 通信用于 KV cache 跨节点共享的实践。

5.6 AMD chiplet 架构催生跨层协同优化 #

MI300X/MI355X 的 8-XCD chiplet 架构在计算效率上引入了 NVIDIA 单 die 不存在的 NUMA 效应。NUMA-Attn 用 ~15 行 Triton 代码解决 attention 的 L2 miss 问题 [2511.02132],HipKittens 在更底层设计了 chiplet-aware scheduling [2511.08083],ConCCL 利用 MI300X 独有的 SDMA 引擎做 DMA offload [2412.14335],Fleet 将 chiplet-task 提升为 programming model 级别的抽象 [ref:L2:2604.15379]。

5.7 集中式 vs 去中心化部署的根本张力 #

Parallax 与 DualPath/PrfaaS/MFS 代表了 LLM 部署架构的两极 [2509.26182] [2602.21548]。MARLIN 占据两者之间的位置——保持 DC 内部的集中式管控,但在 DC 之间用去中心化 multi-agent RL 做 meta-scheduling [2605.13496]。RTP-LLM 的 Master-Worker 架构属于典型集中式设计——全局 Master 做统一 hash map 查找和调度决策 [2605.29639],这在阿里单集群规模下有效,但跨 DC 场景需要与 PrfaaS/GORGO 的跨区域协调机制结合。

5.8 GPU 利用率收割的三种正交维度 #

Harli 收割闲置 SM 算力(~40% → +46% 吞吐)[2511.11729],DualPath 收割闲置 SNIC 带宽 [2602.21548],ZeRO-Prefill 收割 prefill 计算窗口的背景传输能力 [2605.02960]。RTP-LLM 的 cache 感知调度收割了第四种维度——分散在多 worker 上的已计算但未被利用的 KV cache blocks,通过统一 hash map 全局可见并路由复用 [2605.29639]

5.9 CPU 角色定位的四种哲学 #

FastDecode 将 CPU 定位为计算节点(attention 完全在 CPU 执行)[2403.11421],NEO 定位为协处理器(本地 pipelining)[2411.01142],Mooncake 定位为存储池(CPU DRAM 不做计算)[2407.00079]。Blink 走向第四种极端——CPU 仅在启动时加载模型和捕获 CUDA graphs,之后完全退出 steady-state 执行路径 [2604.07609]。RTP-LLM 保留 CPU 作为智能调度中枢——Master 在 CPU 上运行统一 hash map 查找、评分计算和流量调度 [2605.29639],但将高频推理循环(prefill/decode execution)完全交给 GPU。这代表了第五种定位:CPU 做低频高智能决策,GPU 做高频计算,通过降低 CPU 决策频率(20ms 状态轮询 [2605.29639])避免 CPU-Slowdowns 揭示的高频交互瓶颈 [2603.22774]

5.10 RDMA 数据面如何重塑推理系统通信层 #

FaRM 的 RDMA-write circular-buffer messaging 在 2014 年建立了三个关键设计原则 [farm-nsdi14]:(1) 发送端写入接收端预分配 ring buffer(zero-copy);(2) 接收端通过 cache-line versioning 实现 lock-free 一致性读取;(3) PhyCo 2 GB 物理连续页避免 NIC page table thrashing(否则性能 4× 退化)[farm-nsdi14]。十年后这些原则在 LLM serving 中以新形态重现:Blink 的 GPU-resident ring buffer(4096 固定 slot,DPU one-sided RDMA write,GPU atomic CAS 所有权转移)[2604.07609] 是 FaRM ring buffer 从 CPU DRAM 迁移到 GPU memory 的直接继承。OnePiece 的 double-ring buffer(buffer region + size region + CAS spinlock)[2601.20655] 则是 FaRM messaging 从固定大小消息扩展到变长 AIGC tensor 的泛化。KRCore 的 DCT 虚拟化 [2201.11578] 解决了 RDMA ring buffer 方案的一个隐含假设——需要预建连接——通过 5.4μs 连接建立使 RDMA 适用于 serverless 推理的动态连接模式。fabric-lib 的 ImmCounter 原语则证明 "reliable-but-unordered" 是跨 NIC vendor 的 RDMA 最大公约数 [2510.27656]。RTP-LLM 的 RDMA 层(4 层 cache 层级第三层)用 NCCL IBRC 做 KV cache 节点间传输 [2605.29639],代表了 RDMA 在生产 serving 系统中作为 cache 传输通道的成熟应用。

5.11 可编程 serving API vs monolithic 引擎 vs workflow 编排 #

Pie、FlowMesh 和 vLLM/SGLang 代表了 serving 系统设计的三种哲学 [2510.24051] [2510.26913]。Pie 从内部分解生成循环——将 pipeline 拆为可编程 handler。FlowMesh 不改变引擎内部——在 workflow 层做跨租户优化。Blink 代表了第四种哲学——不改变 serving API(保持 OpenAI-compatible),但彻底重构底层执行引擎为 DPU + GPU 双芯片系统 [2604.07609]。RTP-LLM 代表了第五种哲学——保持 OpenAI-compatible API 和标准推理循环,但通过全栈集成(cache 调度 + PD 解离 + 推测解码 + 多模态解耦 + 加载优化)在每个环节累积优势 [2605.29639],而非在单点做革命性重构。

5.12 FlashAttention-4 如何扩展 kernel-framework-hardware 三体关系 #

FlashAttention-4 的设计完全由 Blackwell 的非对称硬件 scaling 驱动——MMA 翻倍但 MUFU exp(16 ops/clock)和 smem 带宽(128 B/clock)不变 [2603.05451]。其三个核心优化均为 Blackwell-specific,不可直接移植到 Hopper 或 AMD CDNA。

5.13 跨区域路由的多层次架构 #

GORGO 和 MARLIN 从不同角度攻击跨区域 LLM 部署。GORGO 用简洁的 additive cost model 将三个异构信号转换为统一的时间量 [2602.11688]。MARLIN 则在更高层次用 game-theoretic multi-agent RL 做 sustainability-aware 跨 DC 调度 [2605.13496]

5.14 RDMA 调度器解耦:从 CPU control plane 到硬件分布式调度 #

CPU-Slowdowns 和 Blink 共同建立了一个新的技术轴:将推理 serving 的控制面从 CPU 迁移到 RDMA fabric + 专用硬件。这一轴与已有的存储/计算/执行/成本维度正交,但与每一维度都有深刻交互。

与存储维度的交互:DualPath 的 CNIC-centric 设计已在数据面使用 RDMA 加载 KV-Cache [2602.21548],但控制面仍在 CPU。Blink 证明控制面本身也可通过 RDMA ring buffer 迁移到非 CPU 硬件 [2604.07609]。组合路径:DualPath 的数据面 + Blink 的控制面 = 全 RDMA 推理栈。

与执行效率维度的交互:GPUOS 和 Blink 的 persistent kernel 共享同一底层技术——GPU-resident 无限运行的 CUDA kernel + lock-free ring buffer 接收外部输入 [2604.17861] [2604.07609]。但目标不同:GPUOS 用 persistent kernel 消除 per-op launch overhead,Blink 用 persistent kernel 消除 per-token host round-trip。Device-side CUDA graph launch(fire-and-forget at ≈2μs,5–8× faster than host launch)[2604.07609] 是两者共享的关键 enabling capability。

与成本维度的交互:CPU-Slowdowns 发现增加 16 vCPU 仅增 ~1.5% 实例成本即可获得 1.36–5.40× TTFT 改善 [2603.22774]——这是 CPU-centric 架构的低成本修补。Blink 需要额外的 BlueField-3 DPU 硬件 [2604.07609]——更高的一次性投入换取完全的 interference immunity。RTP-LLM 的 75% prefill 机器缩减 [2605.29639] 展示了第三条成本路径——通过更智能的调度减少所需机器数量,而非优化单机效率。三条路线的经济学取决于 GPU 利用率和 co-location 密度。

与 RackSched 的 in-network scheduling 的关系:RackSched 在 ToR 可编程交换机上实现了 per-request 动态负载感知调度——LoadTable + ReqTable + tree-based parallel min,全在交换机数据面完成,线速处理 [racksched-osdi20]。其 INT piggyback 机制(reply 包携带服务器队列长度)实现了亚微秒反馈环路 [racksched-osdi20]。将此机制映射到 LLM serving:各 GPU server 的 pending request count / KV-cache 占用 通过 reply 包 piggyback → ToR 交换机做 power-of-k-choices 在 prefill/decode nodes 间负载均衡,零 CPU 调度延迟。RackSched 的 25% Stateful ALU 资源消耗 [racksched-osdi20] 暗示调度逻辑的复杂度受限于 ASIC 资源,LLM 调度的复杂策略(SLO-aware、priority-based)可能超出交换机数据面能力——需要 DPU + switch 的分层协作。RTP-LLM 的 cache-aware 评分函数 [2605.29639] 比 RackSched 的 power-of-k-choices 复杂得多(需要 hash map 查找 + 多因子评分),进一步确认纯 in-network 方案对 LLM 调度的不适用性。

5.15 多模态推理的解耦部署 #

RTP-LLM 的 EPD(Encoder-Prefill-Decode)解离代表了多模态 serving 的新方向——将 ViT 和 LLM 部署在独立 GPU stream 上,造成极端但有意的内存非对称(GPU0 仅 9,280 MB vs GPU1 的 89,088 MB),释放的 GPU0 内存允许更高 ViT batch 并发,实现 1.86–2.52× 多模态吞吐提升 [2605.29639] [2605.29639]。这一设计本质上是 PD 解离思想在多模态维度的推广——不仅 prefill 和 decode 物理分离,encoder 和 decoder 也物理分离。

5.16 FlashAttention 谱系:attention kernel 的硬件代际协同演进 #

FlashAttention 系列构成了 §1 执行效率维度中 FA4 的完整前史,也是本 topic「attention kernel 随硬件重新设计」方法论的源头。FA1 的核心洞察是 attention 本质 memory-bound——通过 tiling + online softmax + recomputation 避免物化 $N^2$ 中间矩阵,把 HBM I/O 从 $O(N^2)$ 降到 $O(N)$ [2205.14135]。这一 IO-aware 视角与 §1 计算维度的 Ops:Byte 分析 [gpu-arch-comparison-2025] 是同一枚硬币的两面——都指向"数据搬运而非算力是瓶颈"。FA2 发现 FA1 的 GPU 利用率仅 25–40%,根因是非 matmul FLOPs 过多、并行维度选择次优、warp 间通信冗余——重划分工作后达 50–73% peak [2307.08691]。FA3 则进入 Hopper 的异步时代——warp specialization(producer/consumer warp 分工)+ TMA/WGMMA 异步流水线 + FP8 低精度,将 attention 从"同步 tiling"推向"异步 pipeline"[2407.08608],其 warp specialization 思想与 DeepSeek-V3 DualPipe 的通信 warp 专门化 [2412.19437] 属于同一硬件趋势的两个应用面。FA1→FA2→FA3→FA4 [2603.05451] 每一代都对应一次硬件代际(Ampere→Ampere 优化→Hopper→Blackwell),完美印证了 §5.12 提出的"kernel-framework-hardware 三体关系"——attention kernel 不再是可移植的通用实现,而是硬件代际专用设计。同时,NUMA-Attn 揭示的 FlashAttention2 在 MI300X chiplet 上 L2 命中率仅 ~1% [2511.02132] 说明 FA 谱系的 NVIDIA-first 设计在 AMD chiplet 上需要重新映射,进一步强化跨 vendor 不可移植性。

5.17 推测解码从算法单点到系统级调度的演进 #

推测解码在本 topic 呈现出与执行效率维度深度耦合的系统化趋势。早期推测解码是纯算法问题(draft-then-verify),但新一波工作把它变成调度问题。SSD 对推测过程递归地再推测,压缩草稿本身的开销 [2603.03251]。SPECTRE 认识到并行推测在高负载下会争抢目标模型批容量,故混合普通解码与并行推测以在资源受限下最优 [2605.08151]。DSpark 把这一系统视角推到极致——它明确把"选多长验证"formulate 成全局吞吐最大化 $\Theta=\tau\cdot\text{SPS}(B)$,按实测 SPS(B) 曲线动态调整每请求验证预算:中等并发时用空闲目标算力扩到 4–6 token,高并发时收缩剪掉低置信尾 token 保批容量 [dspark]。这与 §1 成本维度的核心矛盾(SLO 达标率 vs 吞吐的 Pareto 前沿 [2504.09285])直接呼应——DSpark 把推测长度变成一个可随负载调节的旋钮来外推 Pareto 前沿。生产化层面,DSpark 与 RTP-LLM 的模块化 C++ 推测解码框架 [2605.29639] 以及 DeepSeek-V3 MTP [2412.19437] 构成了从训练侧(产出 MTP/draft heads)到服务侧(负载自适应验证调度)的端到端闭环——推测解码已从"更快解码的技巧"演变为"serving 系统的一等调度维度"。

5.18 KV cache 复用的粒度扩张:从前缀到跨模型到多 agent 集体共享 #

KV cache 复用的演进呈现出一条清晰的粒度扩张主线。最初的复用发生在单请求前缀内(PagedAttention 的 block 级 CoW [2309.06180]),SGLang 的 RadixAttention 把它扩展到跨请求前缀共享 [2312.07104],Mooncake 进一步扩展到跨节点分布式前缀池 [2407.00079]。新一波工作突破了"前缀"和"单模型"两个边界:Marconi 处理 hybrid 模型下 SSM 状态无法按前缀切分的问题,把复用扩展到非纯 attention 架构 [2411.19379];LMCache 把复用位置从"仅前缀"扩展到"任意位置的 KV 片段" [2510.09665];ICaRus 跨越模型边界,发现不同模型间存在可复用的 identical cache [2603.13281];TokenDance 则在 agent 维度实现 collective KV cache sharing——同一 workflow 内多个 agent 集体共享上下文 KV [2604.03143]。这条粒度扩张主线与 §5.1(存储约束计算)和 §6 共识 7(cache-centric scheduling)深度交织——KV cache 正从"内存管理对象"演变为"跨请求/跨模型/跨 agent 的一等共享资源"。CacheGen 从编码角度补齐了这一图景——当复用发生在带宽受限的跨节点/跨层级路径上时,KV cache 需要像视频流一样被压缩编码并带宽自适应地流式加载 [2310.07240]。KVDrive 和 Resident KV Claims 则从系统契约角度形式化了多层复用的生命周期管理 [2605.18071] [2605.24259],与 RTP-LLM 的四层 cache 层级 [2605.29639] 属同一多层管理范式。

5.19 MoE serving 的三层优化栈:从通信到 dispatch 到 decode 路由 #

MoE serving 优化在本轮呈现出清晰的三层分工。最底层是通信优化(§5.3)——DeepSeek-V3 DualPipe 隐藏 all-to-all [2412.19437]、ZeRO-Prefill 消除 all-to-all [2605.02960]、MFS 隔离网络竞争 [2603.17456]。中间层是 dispatch-time 负载均衡——SGLang 的 Waterfill 把 dense shared expert 当可分派 slot 填谷到轻负载 rank,LPLB 用每层 min–max LP 在 redundant 副本间实时重分流量,两者都作用于"物理落点"而不改 logical top-k,是可放进推理关键路径的低开销运行时机制(三次 CUDA kernel launch)[blog-waterfill-lplb]。最上层是 decode 路由——ELDR 发现 PD-disaggregated MoE 下 decode 实例选择应感知 expert 放置局部性,减少跨节点 expert 访问 [2607.00466];NetKV 从网络代价角度做 decode instance selection [2606.03910]。这三层与 §1 成本维度和 §5.3 模型架构约束共同构成 MoE 生产部署的完整优化栈——通信层降绝对开销、dispatch 层平单 batch 残余不均衡、路由层减跨节点访问,三者正交互补。

§6 共识与分歧 #

共识 #

  1. PagedAttention 是事实标准 KV cache 管理机制:vLLM 的 block table + CoW 分页方案 [2309.06180] 被几乎所有后续系统继承——Mooncake 扩展为分布式 hash-indexed block pool [2407.00079],Blink 在 GPU-resident ring buffer 中保留了 paged KV-cache 管理(完全在 GPU 执行)[2604.07609],RTP-LLM 在其上构建了带 watermark 的采样前缀哈希——满块使用引用计数允许多请求并发访问,部分块为独占模式 [2605.29639]
    1. PD 分离是必要的基础架构:所有涉及推理 serving 的系统均建立在 prefill 与 decode 分离或至少 decode-first 优化的前提上。DualPath [2602.21548]、PrfaaS [2604.15039]、DynaServe [2504.09285]、PPD [2603.13358]、KVServe [kvserve]、RTP-LLM [2605.29639] 六篇直接在 PD 架构上创新。RTP-LLM 进一步验证了 PD 解离在 480B MoE 规模的生产可行性——4.72–5.33× TTFT 加速且吞吐不变 [2605.29639]
      1. RDMA 是大规模 GPU 集群数据搬运的必需原语:DualPath [2602.21548]、TensorHub [2604.09107]、fabric-lib [2510.27656]、Blink [2604.07609]、OnePiece [2601.20655]、RTP-LLM [2605.29639] 六篇独立工作均选择 RDMA 作为核心传输机制。RTP-LLM 将 RDMA 用于 4 层 KV cache 层级的第三层(远程 CPU via RDMA)和 PD 解离的 KV cache 传输(NCCL IBRC)[2605.29639]。FaRM 十年前的 ring buffer messaging 范式 [farm-nsdi14] 在 Blink 和 OnePiece 中得到直接继承——证明 RDMA 编程模型具有跨代际的普适性。
        1. CPU 位于推理 critical path 是架构级脆弱性,而非资源不足:CPU-Slowdowns 证明即便 CPU 核心充足,broadcast 竞争和 barrier 放大仍是结构性问题 [2603.22774]。Blink 在 isolation 条件下(无 CPU 干扰)仍获得最高 8.47× P99 TTFT 改善,证明 CPU 开销不仅是 interference 问题更是 baseline overhead [2604.07609]。GPUOS 的 persistent kernel 从另一角度验证——消除 CPU per-op dispatch 获得 15.3× micro-op 加速 [2604.17861]。RTP-LLM 通过降低 CPU 决策频率(20ms 轮询)部分缓解了此问题 [2605.29639],但其 Master 仍运行在 CPU 上。
          1. Agentic workload 正在重塑 serving 系统假设:至少 8 篇工作以 agentic 场景为核心评估——KVFlow [2507.07400]、DualPath [2602.21548]、ThunderAgent [ref:L2:2602.13692]、Helium [ref:L2:2603.16104]、Halo [ref:L2:2509.02121]、Concur [ref:L2:2601.22705]、Sutradhara [2601.12967]、GLM Graph-CoT [2511.01633]、TokenDance [2604.03143] 等。其共同前提——跨请求/跨 agent 前缀复用——的技术源头是 SGLang 的 RadixAttention [2312.07104]
            1. Inter-kernel overhead 在低 batch size 下是性能主要瓶颈:Three Taxes [2511.02168]、TileRT [tilert-speed-scaling-law]、GPUOS [2604.17861]、Blink [2604.07609] 四篇从不同角度验证。Blink 特别揭示了 MoE 模型的放大效应:Qwen-3 30B-A3B 仅激活 3B/30B 参数使每步计算极快,但 CPU orchestration cost 不变——调度 overhead 占步骤时间的比例更大 [2604.07609]
              1. Cache-centric scheduling 正成为 PD 解离系统的核心范式:Mooncake 最先明确提出 "KVCache-centric" 概念 [2407.00079],RTP-LLM 用更大规模生产数据验证——统一 hash map + cache 感知评分函数 + 4 层持久化层级,215% cache 复用提升和 75% prefill 机器缩减 [2605.29639] [2605.29639]。PrfaaS 将此范式扩展到跨 DC [2604.15039]。LMCache 把 cache 层抽象为独立可插拔中间件进一步固化此范式 [2510.09665],云 trace 分析(KVCache in the Wild)为 workload-aware 驱逐提供实证基础 [2506.02634]。KV cache 正从"内存管理对象"演变为"一等公民调度资源"。
                1. Attention kernel 已从通用实现转向硬件代际专用设计:FlashAttention 完整谱系(FA1 IO-aware tiling → FA2 work partitioning → FA3 Hopper 异步/warp specialization → FA4 Blackwell TMEM)证明每次硬件代际都需要重新设计 attention 数据流 [2205.14135] [2307.08691] [2407.08608] [2603.05451],NUMA-Attn 在 AMD chiplet 上的 L2 重映射 [2511.02132] 与 MPK 的通用 mega-kernel 编译 [2512.22219] 从跨 vendor 与自动化两个方向延续此趋势。
                  1. 推测解码正在成为 serving 的一等调度维度而非纯算法技巧:SSD [2603.03251]、SPECTRE [2605.08151]、DSpark [dspark] 与 RTP-LLM 的模块化框架 [2605.29639] 共同验证——验证长度/草稿深度应随系统负载动态调节,DeepSeek-V3 MTP [2412.19437] 提供训练侧闭环。
                  2. 分歧 #

                    1. CPU 角色定位的五元分歧。FastDecode 将 CPU 定位为计算节点 [2403.11421]。NEO 定位为协处理器 [2411.01142]。Mooncake 定位为存储池 [2407.00079]。Blink 彻底消除 CPU 从 steady-state [2604.07609]。RTP-LLM 保留 CPU 为低频智能调度中枢(20ms 轮询周期,仅做路由决策不参与推理循环)[2605.29639]。CPU-Slowdowns 的数据为这一分歧提供了量化依据:5 核 → 32 核约 9× 加速 [2603.22774] 证明 CPU 加核心有效但边际递减。RTP-LLM 的路线暗示:如果 CPU 仅做低频全局决策(而非高频 per-token 调度),CPU 瓶颈可被架构性规避而无需完全移除。
                      1. RDMA ring buffer 的 safety-liveness trade-off。FaRM 的 lock-free reads 通过 cache-line versioning 实现了严格可串行化一致性——正确性依赖 x86 DMA cache coherence + RDMA increasing-address-order write 两个硬件保证 [farm-nsdi14]。OnePiece 的 double-ring buffer 在仅有 CAS 原子操作可用的约束下,选择了放松 safety 换取 liveness [2601.20655]。Blink 的 ring buffer 走中间路线——per-slot state machine + atomic CAS 所有权转移 [2604.07609]。三者在 safety-liveness-performance 三角中占据不同位置。
                        1. 调度硬件化程度的谱系。RackSched 将调度逻辑完全下沉到 ToR 交换机 ASIC [racksched-osdi20]。Blink 在 DPU(16 ARM 核心)执行 request 管理 [2604.07609]。RTP-LLM 在主机 CPU 上运行 cache-aware 评分函数 [2605.29639]。传统方案(vLLM/SGLang)在主机 CPU 执行全部调度。四者的 trade-off 轴是"策略复杂度 vs 执行延迟 vs interference immunity"——RTP-LLM 选择了最高策略复杂度(多因子评分 + hash map 查找),以牺牲部分 interference immunity 换取更优调度质量。
                          1. RDMA 连接建立成本是否仍是瓶颈。KRCore 证明标准 RDMA verbs 的连接建立(15.7ms/QP)在 elastic 场景下是致命瓶颈,DCT 虚拟化将其降至 5.4μs [2201.11578]。但 Blink、OnePiece 和 RTP-LLM 均假设 RDMA 连接在启动时预建——RTP-LLM 的 RDMA 层用于固定节点间 KV cache 传输 [2605.29639]。KRCore 的价值在 LLM 场景中主要体现在 serverless cold-start 和动态 scaling。
                            1. 存储层方案选择:DRAM pool vs SSD + NIC 聚合 vs 模型压缩 vs KV 压缩 vs 多层层级。Mooncake 用分布式 DRAM pool [2407.00079],DualPath 选择 SSD 存储 + NIC 带宽聚合 [2602.21548],PrfaaS 走 hybrid attention 压缩 [2604.15039],KVServe 走 service-aware 在线压缩 [kvserve],RTP-LLM 走四层 cache 层级 + 自适应量化(FP16→INT8/INT4/FP8)[2605.29639]。RTP-LLM 的方案独特之处在于不依赖单一机制而是组合——每一层承担不同容量/延迟角色 [2605.29639]
                              1. MoE 通信:消除 AllToAll vs 优化 AllToAll vs 调度 AllToAll。ZeRO-Prefill 完全消除 [2605.02960],DualPipe 极致优化 [2412.19437],MFS 隔离竞争 [2603.17456]。RTP-LLM 选择直接使用 DeepEP 做 All2All [2605.29639]——采纳而非重新发明通信优化。
                                1. 通信融合粒度的五级分化。Three Taxes 在 kernel 内嵌入 tile-level pipeline [2511.02168];DualPipe 在 framework 层实现组件级重叠 [2412.19437];TileRT AOT 编译全模型 [tile-ai-tilert];GPUOS 用 persistent kernel + JIT [2604.17861];Blink 用 persistent kernel + RDMA ring buffer 彻底消除 host 参与 [2604.07609]
                                  1. 全栈集成 vs 单点突破的系统设计哲学。RTP-LLM 走全栈集成路线——将 PD 解离、cache 调度、推测解码、多模态解耦、加载优化集成到同一系统 [2605.29639],每个组件的增益中等(1.12× 推测解码、35–37% TTFT、6.27× 加载)但组合效果显著。Blink 走单点突破——仅攻击 CPU 在 critical path 中的角色但获得 8.47× P99 TTFT [2604.07609]。TileRT 同样是单点极致——仅攻击 kernel boundary overhead [tile-ai-tilert]。两种哲学的 trade-off:全栈集成提供更广泛的场景覆盖但单个组件贡献难以消融 [2605.29639];单点突破提供更清晰的贡献归因但场景适用性窄。
                                    1. KV cache 复用粒度的分歧:前缀 vs 任意位置 vs 跨模型 vs 跨 agent。SGLang/Mooncake 复用共享前缀 [2312.07104] [2407.00079],LMCache 复用任意位置的 KV 片段 [2510.09665],Marconi 处理 hybrid 模型下非前缀切分的 SSM 状态 [2411.19379],ICaRus 跨模型复用 identical cache [2603.13281],TokenDance 在多 agent 间集体共享 [2604.03143]。粒度越粗,复用机会越多但一致性/命中判定越复杂——尚无统一框架涵盖全部粒度。
                                      1. CPU-GPU 协同的时机分歧:offload 计算 vs 重叠执行 vs 完全移除。FastDecode/NEO 把计算 offload 到 CPU [2403.11421] [2411.01142],APEX 走异步并行重叠路线——CPU 与 GPU 阶段流水重叠而非各自承担独立计算 [2506.03296],Blink 则完全移除 CPU 从 steady-state [2604.07609]。三者对"CPU 在推理中该扮演什么角色"给出不同答案,取决于目标硬件(资源受限设备偏向 offload/overlap,数据中心偏向移除)。
                                        1. PD 传输/decode 选择的优化信号分歧。FlowKV 优化 KV 传输通道本身的延迟 [2504.03775],NetKV 以网络代价为信号选 decode 实例 [2606.03910],ELDR 以 expert 局部性为信号做 MoE decode 路由 [2607.00466],RTP-LLM 以 cache 局部性为信号 [2605.29639]。四者选择了不同的主导信号(传输延迟 / 网络拓扑 / expert 放置 / cache 命中),反映 PD 解离下 decode 端调度尚无收敛的单一最优目标。
                                        2. §7 Open challenges (根本性困难) #

                                          7.1 KV-Cache 压缩与 attention 表达力的信息论瓶颈 #

                                          Hybrid attention 可将 KV 吞吐降低 4–13× [2604.15039],KVServe 可达 10× [kvserve],RTP-LLM 的自适应量化(FP16→INT8/INT4/FP8)提供运行时压缩 [2605.29639],但缺乏理论工具量化压缩极限。RTP-LLM 的四层层级延长了 cache 生命周期但未减少 cache 数据量——当所有四层都满时仍需驱逐 [2605.29639]。预估难度:3–5 年。

                                          7.2 异构硬件下的通用性能可移植 #

                                          DualPath 绑定 InfiniBand VL QoS [2602.21548];Blink 绑定 NVIDIA BlueField-3 DPU + ConnectX-6 [2604.07609];KRCore 绑定 Mellanox ConnectX 系列的 DCT 私有扩展 [2201.11578];TileRT 绑定 NVLink TMA [tile-ai-tilert];Fleet 绑定 MI350 chiplet-task [ref:L2:2604.15379];RTP-LLM 绑定阿里内部 3FS 分布式存储和 FUSE 文件系统 [2605.29639]。fabric-lib 尝试统一 ConnectX-7 和 EFA [2510.27656],但 NIC vendor 专有功能使跨 vendor 可移植性停留在最大公约数。预估难度:2–3 年。

                                          7.3 动态 agentic workflow 的在线资源分配 #

                                          KVFlow 假设 workflow graph 运行时稳定可观测 [2507.07400],但 ReAct/Reflexion 等自主 agent 在运行时动态生成下一步。Concur 的 AIMD 是最少假设的方案 [ref:L2:2601.22705]——但无法利用 workflow 结构信息做预取。RTP-LLM 的 chat-ID 亲和路由 [2605.29639] 是 agentic 场景的实用方案(将同一对话的后续轮次路由到同一 worker 复用 KV cache),但无法处理跨 chat 的动态分支。预估难度:1–3 年。

                                          7.4 千 GPU 级容错与一致性 #

                                          DualPath 部署于 1152 GPU [2602.21548],TensorHub 部署于 1024 GPU [2604.09107]——但容错机制均未充分披露。RTP-LLM 部署规模覆盖 100M+ 用户 [2605.29639],其 Name Service + Carbon 自动故障恢复 [2605.29639] 是目前披露最完整的生产容错方案之一,但论文未详细描述 cache 一致性在节点故障时的行为。RDMA ring buffer 的故障恢复尤其困难:one-sided operations 的远端无 CPU 参与意味着远端 crash 时无法主动通知发送方 [farm-nsdi14]。预估难度:3–5 年。

                                          7.5 AOT 静态编译 vs 动态工作负载的根本性张力 #

                                          TileRT 的 AOT 实现了 inter-kernel overhead 的彻底消除 [tile-ai-tilert],但现代 LLM 的动态性不断增加。Blink 的 persistent scheduler 面临类似张力——CUDA graph cache 需要预编译 650–1000 个 (batch, seq_len) 对的图 [2604.07609],unseen configurations 需要 fallback。RTP-LLM 的模块化推测解码同样受限于此——step size = 1 的保守配置 [2605.29639] 暗示更大 step size 下的动态分支处理尚未成熟。预估难度:2–4 年。

                                          7.6 MoE 通信范式的根本选择尚未收敛 #

                                          ZeRO-Prefill 引入 168 GB/layer 的 AllGather [2605.02960];DualPipe 的 warp specialization 将 15% SM 永久分配给通信 [2412.19437];MFS 的 RMLQ 需要精确 flow 分类 [2603.17456]。RTP-LLM 直接使用 DeepEP [2605.29639] 回避了这一设计决策但继承了其局限。预估难度:2–3 年。

                                          7.7 跨节点 in-kernel fusion 的空白地带 #

                                          fabric-lib 占据 inter-node + library-level + async P2P 位置,Three Taxes/Iris 占据 intra-node + in-kernel + fine-grained sync 位置,交集区域完全空白。预估难度:1–2 年。

                                          7.8 CPU-free 推理的多 GPU 扩展 #

                                          Blink 当前限于单 GPU [2604.07609]。多 GPU 扩展需要解决三个根本问题。RTP-LLM 的 CPU-based Master 在多 GPU 场景下不受此限制(Master 本身不在推理 critical path,仅做路由决策)[2605.29639],但牺牲了 Blink 的 interference immunity。预估难度:2–4 年。

                                          7.9 Persistent kernel + dynamic batching 的融合 #

                                          TileRT/TokenSpeed 局限于 BS=1 [tile-ai-tilert] [tokenspeed]。Blink 实现了 persistent scheduler + continuous batching 的组合 [2604.07609],但其 FCFS 策略过于简单。RTP-LLM 的 cache-aware 评分函数 [2605.29639] 展示了生产级调度所需的复杂度(多因子评分 + 预测调度 + chat-ID 亲和),但这些决策目前运行在 CPU 上——将此复杂度嵌入 GPU-resident scheduler 是开放挑战。预估难度:1–3 年。

                                          7.10 Cache-aware 调度的可复现性与负载依赖性 #

                                          RTP-LLM 的评分函数权重 $\alpha, \beta, \gamma$ "基于负载特征调优"但未公布数值 [2605.29639],生产结果显示极端的负载依赖性——机器人 Q&A 获得 215% cache 复用提升而淘宝商家仅 <1% [2605.29639] [2605.29639]。这暗示 cache-centric scheduling 的收益高度依赖 workload 特征(prefix 共享度、query 多样性),缺乏通用的"最优调度策略"——每个部署场景需要独立调优。预估难度:持续性工程挑战。

                                          §8 成熟度判断 #

                                          技术方向成熟度证据趋势
                                          PD 分离推理Production-readyDualPath 1152 GPU 生产 [2602.21548];PrfaaS Moonshot AI 生产 [2604.15039];RTP-LLM 100M+ 用户生产(阿里)[2605.29639]加速
                                          Cache-centric SchedulingProduction-readyRTP-LLM 全栈集成生产验证——75% prefill 机器缩减、45% cache 命中率 [2605.29639];Mooncake KVCache-centric 架构 [2407.00079]加速
                                          RDMA P2P 通信Production-readyfabric-lib 部署于 Perplexity AI [2510.27656];TensorHub ByteDance RL [2604.09107];RTP-LLM 用 NCCL IBRC 做节点间 KV 传输 [2605.29639];FaRM Microsoft 内部长期运行 [farm-nsdi14]加速
                                          Persistent Kernel / AOT 引擎Early productionTileRT 驱动 GLM-5.1 生产 [tile-ai-tilert];TokenSpeed 驱动 Kimi K2.5 生产 [tokenspeed];Blink H100 + BlueField-3 完整实现 [2604.07609]加速
                                          推测解码生产化Production-readyRTP-LLM 模块化 C++ 框架支持 4 种算法,DeepSeek-V3 生产部署 ~1.9 tokens/step 接受率 [2605.29639]加速
                                          多模态解耦部署Early productionRTP-LLM EPD 架构生产验证——1.86–2.52× 多模态吞吐 [2605.29639]加速
                                          快速模型加载Production-readyRTP-LLM 文件序驱动加载 6.27× 加速,生产环境分钟级 600B+ 模型更新 [2605.29639]稳定
                                          CPU-free inferenceEarly researchBlink 单 GPU 验证 4 模型 × 3 baselines [2604.07609];需 DPU 硬件;单 GPU 限制;无多 GPU 支持加速
                                          RDMA 控制面加速Research-frontierKRCore 开源验证 5.4μs 连接 [2201.11578];OnePiece 无实验验证 [2601.20655]加速
                                          In-network μs schedulingConcept → Early research (for LLM)RackSched 12-server testbed 验证 μs 级 RPC 调度 [racksched-osdi20];无 LLM-specific 验证观察
                                          SLO-aware 调度Research-frontierJITServe 1/8.55 竞争比 [2504.20068];Justitia KV token-time [ref:L2:2510.17015];MFS RMLQ [2603.17456];均未大规模生产部署加速
                                          MoE 执行优化Research-frontier → Early productionZeRO-Prefill vLLM v0.11 开源 [2605.02960];RTP-LLM 集成 DeepEP 部署 480B MoE [2605.29639]加速
                                          Agentic workflow 服务Early research → Research-frontierThunderAgent 3.58× [ref:L2:2602.13692]、Halo 400× [ref:L2:2509.02121]、Concur 4.09× [ref:L2:2601.22705]——但均未千卡验证快速加速
                                          RL 训练基础设施Production-readyveRL EuroSys 2025 [2409.19256];TensorHub 1024 GPU [2604.09107]稳定
                                          AMD 软件栈效率Migrating从 45% 利用率向 93% 追赶 [2510.27583];HipKittens BF16 64.4% peak [2511.08083]加速

                                          整体判断:本轮新增 RTP-LLM 的最大信号是验证 cache-centric scheduling 已达 production-ready 成熟度。此前 Mooncake 提出 KVCache-centric 概念但规模有限;RTP-LLM 在 100M+ 用户、480B MoE 模型规模上全面验证了 PD 解离 + 多层 KV cache + cache-aware 路由的生产可行性 [2605.29639]。同时,RTP-LLM 的全栈集成设计(模块化推测解码 + EPD 多模态解耦 + 文件序加载)确立了"生产级推理引擎"作为独立系统类别——区别于学术单点突破(Blink、TileRT)和通用开源框架(vLLM、SGLang)[2605.29639]

                                          §9 邻接 topic #

                                          1. Agent Serving(agent 运行时优化):与本 topic 在 agentic workload characterization 上深度重叠。边界在于本 topic 聚焦"集群级资源管理",agent-serving 更关注"agent 运行时的推理引擎优化"。
                                            1. RDMA 通信(跨硬件可移植通信):fabric-lib、TensorHub、FaRM、KRCore、OnePiece、RTP-LLM(RDMA 层)的 RDMA 层工作同时属于本 topic 和独立通信 topic。本轮 cluster 类别已达 8 篇,RDMA 通信可能需要独立 topic。
                                              1. MoE 推理优化:DeepSeek-V3、ZeRO-Prefill、MFS、FaaSMoE 以及 RTP-LLM 的 480B MoE 部署涉及 MoE 特有挑战。
                                                1. GPU 架构演进:MI300X/Blackwell/CDNA4 白皮书四篇提供硬件基础约束。
                                                  1. Persistent Kernel 执行范式:TileRT、TokenSpeed、Fleet、GPUOS、Blink 构成了从 per-op JIT 到 full-model AOT 到 full-system persistent scheduler 的完整频谱。
                                                    1. SmartNIC/DPU 卸载:Blink 是目前 LLM serving 中唯一完整利用 DPU 的系统 [2604.07609]
                                                      1. In-network Computing:RackSched [racksched-osdi20] 代表了将调度逻辑下沉到网络硬件的极端路线。
                                                        1. Geo-Distributed & Sustainability-Aware Scheduling:MARLIN [2605.13496] + GORGO [2602.11688] + PrfaaS [2604.15039] 构成跨区域 serving 完整栈。
                                                          1. Decentralized/Volunteer GPU Inference:Parallax [2509.26182] 是唯一面向 volunteer WAN pool 的系统。
                                                            1. Inference-Training Co-location:Harli [2511.11729] 开辟了在同一 GPU 上共存 inference+training 的新维度。
                                                              1. 推测解码优化:RTP-LLM 的模块化框架 [2605.29639] 将推测解码从算法研究推进到生产系统集成阶段,与 DeepSeek-V3 MTP [2412.19437] 构成训练-推理端到端闭环。该方向可能独立成 topic。
                                                                1. 多模态 Serving 架构:RTP-LLM 的 EPD 解耦 [2605.29639] 开辟了 ViT-LLM 物理分离的新设计维度。随着多模态模型规模扩大,独立的"多模态推理系统设计"topic 可能出现。
                                                                2. §10 参考 #

                                                                  EntityCategoriesRole in topicKey contribution
                                                                  [farm-nsdi14]clusterRDMA 数据面奠基RDMA-write circular buffer messaging + lock-free one-sided reads,10× 吞吐 / 100× 低延迟,建立 RDMA ring buffer 范式
                                                                  [2201.11578]clusterRDMA 弹性连接DCT 虚拟化将连接建立从 15.7ms 降至 5.4μs(2,900×),O(1) 内存,使 RDMA 适用于 elastic/serverless
                                                                  [racksched-osdi20]cluster硬件化调度范式ToR 可编程交换机 μs 级 per-request 调度,power-of-k-choices,近线性扩展
                                                                  [2601.20655]clusterRDMA 推理通信One-sided RDMA + double-ring buffer 实现 AIGC 多阶段零 CPU 通信
                                                                  [2604.07609]frameworkCPU-free 推理SmartNIC frontend + GPU-resident persistent scheduler,P99 TTFT 降低 8.47×
                                                                  [2603.22774]clusterCPU 瓶颈动机首次系统性量化 multi-GPU CPU 隐藏瓶颈
                                                                  [2309.06180]framework存储奠基PagedAttention: OS 分页 KV cache,block table + CoW,吞吐 2–4×
                                                                  [2407.00079]framework存储+架构奠基Mooncake: KVCache-centric disaggregated 架构,分布式 DRAM pool
                                                                  [2405.04434]framework模型驱动部署DeepSeek-V2 MLA: KVCache 压缩 93.3%
                                                                  [2305.05920]framework调度奠基FastServe: skip-join MLFQ 消除 HoL blocking
                                                                  [2403.11421]frameworkCPU offloadingFastDecode: attention 卸载到分布式 CPU
                                                                  [2411.01142]frameworkCPU offloadingNEO: asymmetric GPU-CPU pipelining
                                                                  [2510.24051]framework可编程 servingPie: Wasm inferlets(SOSP 2025)
                                                                  [2510.26913]frameworkworkflow 编排FlowMesh: DAG identity + utility scheduling
                                                                  [2604.06370]frameworkmulti-LoRA KVForkKV: DualRadixTree
                                                                  [2602.11688]framework跨区域路由GORGO: additive cost model
                                                                  [2603.05451]kernel注意力 kernelFlashAttention-4: B200 71% peak
                                                                  [2602.21548]framework, cluster存储核心聚合 DE 空闲 SNIC,1152 GPU
                                                                  [2604.15039]framework存储+成本跨 DC PD 分离,KV 吞吐降至 3 Gbps
                                                                  [2604.09107]framework, cluster存储核心ROS 无所有权存储 + pipeline replication
                                                                  [kvserve]framework存储Service-aware KV 压缩
                                                                  [2603.13358]framework存储+调度Append-prefill 仅 2% 干扰
                                                                  [2510.27583]hardware计算核心MI300X 45% 利用率诊断
                                                                  [2512.02189]hardware计算核心Blackwell TMEM/tcgen05/FP4 微架构量化
                                                                  [gpu-arch-comparison-2025]hardware计算+成本AMD vs NVIDIA 四代架构对比
                                                                  [amd-cdna4-whitepaper]hardware计算MI355X: FP8 5PF, MXFP4 10PF, 288 GB HBM3E
                                                                  [2511.02168]cluster, kernel计算Three Taxes + Iris compute-comm fusion
                                                                  [2510.27656]cluster计算+存储可移植 RDMA P2P 库,统一 ConnectX-7/EFA
                                                                  [2412.14335]kernel计算DMA offload,C3 加速比 21%→72%
                                                                  [2511.02132]kernel计算Swizzled mapping,MI300X L2 90–97%
                                                                  [2511.08083]kernel计算AMD tile DSL,BF16 1610 TFLOPS
                                                                  [2604.17861]kernel执行效率Persistent kernel + JIT,micro-ops 15.3×
                                                                  [2412.19437]framework, kernel计算核心DualPipe + custom all-to-all
                                                                  [2504.09285]framework成本核心微请求抽象,容量 1.15–3.07×
                                                                  [2504.20068]framework成本核心多类型 SLO + 1/8.55 竞争比
                                                                  [2603.17456]framework成本+MoERMLQ 三阶段流量调度
                                                                  [2605.02960]framework成本+MoEAsyncEP 消除 AllToAll
                                                                  [2605.13496]cluster成本+sustainabilityGame-theoretic multi-agent RL 跨 DC 路由
                                                                  [tile-ai-tilert]framework执行效率AOT Engine Kernel,GLM-5.1 生产
                                                                  [tokenspeed]framework执行效率C++ FSM + Placement,Kimi K2.5 生产
                                                                  [2409.19256]frameworkRL 训练veRL 开源,1.53–20.57× RLHF 加速
                                                                  [2511.11729]frameworkco-locationInference+PEFT co-location
                                                                  [2509.26182]framework去中心化Decentralized volunteer WAN 推理
                                                                  [2604.26881]frameworkMoE serverlessExpert-as-FaaS 多租户共享
                                                                  [2605.29639]framework生产级全栈推理引擎PD 解离 + 4 层 KV cache + cache-aware 调度 + 推测解码 + EPD 多模态 + 文件序加载,100M+ 用户生产验证,480B MoE 4.72–5.33× TTFT,6.27× 加载加速
                                                                  [2205.14135]kernel注意力 kernel 奠基FlashAttention: IO-aware tiling + online softmax,HBM $O(N^2)\to O(N)$
                                                                  [2307.08691]kernel注意力 kernelFlashAttention-2: work partitioning,50–73% peak
                                                                  [2407.08608]kernel注意力 kernelFlashAttention-3: warp specialization + asynchrony + FP8,Hopper 75% peak
                                                                  [2512.22219]kernel执行效率MPK: 通用 mega-kernel 编译器 + 运行时
                                                                  [2603.24517]kernelkernel 自动化AVO: agentic 进化搜索生成/优化 GPU kernel
                                                                  [2310.07240]framework存储/KV 编码CacheGen: KV 算术编码 + 带宽自适应流式加载
                                                                  [2312.07104]framework存储/前缀复用奠基SGLang: RadixAttention 跨请求前缀复用 + FSM 约束解码
                                                                  [2411.19379]frameworkhybrid 模型前缀缓存Marconi: SSM 状态感知的 speculative insertion + FLOP-aware 驱逐
                                                                  [2412.19442]frameworkKV 管理综述KV Cache Survey: token/model/system 三层分类法
                                                                  [2504.03775]frameworkPD KV 传输FlowKV: 低延迟 KV cache 传输框架
                                                                  [2506.02634]framework生产 cache 画像KVCache in the Wild: 云厂商 KV 访问模式 + workload-aware 驱逐
                                                                  [2506.03296]frameworkCPU-GPU 重叠APEX: 资源受限设备异步并行 CPU-GPU 执行
                                                                  [2510.09665]framework企业级 KV cache 层LMCache: 跨引擎/请求/存储层级复用任意位置 KV
                                                                  [2511.01633]frameworkagentic 推理GLM Graph-CoT: 多 agent + 高效 serving 规模化图链推理
                                                                  [2601.12967]frameworkagentic 编排Sutradhara: orchestrator-engine co-design
                                                                  [2603.03251]framework推测解码SSD: 对推测过程再推测
                                                                  [2603.13281]framework跨模型 KV 复用ICaRus: 多模型间 identical cache reuse
                                                                  [2604.03143]framework多 agent KV 共享TokenDance: collective KV cache sharing
                                                                  [2605.08151]framework推测解码SPECTRE: hybrid ordinary-parallel speculative serving
                                                                  [2605.18071]framework多层 KV 管理KVDrive: 长上下文 GPU/CPU/SSD 整体式 KV 管理
                                                                  [2605.24259]frameworkKV 复用契约Resident KV Claims: future-reuse conformance contract
                                                                  [2607.00466]frameworkMoE decode 路由ELDR: expert-locality-aware decode routing
                                                                  [2606.03910]clusterPD decode 选择NetKV: network-aware decode instance selection
                                                                  [dspark]framework推测解码调度DSpark: 半自回归草稿 + 置信度调度验证,无损,Pareto 前沿外推
                                                                  [blog-waterfill-lplb]framework (blog)MoE dispatch LBWaterfill + LPLB: dispatch-time MoE 负载均衡,不改 logical top-k