TraCT: Disaggregated LLM Serving with CXL Shared Memory KV Cache at Rack-Scale

发表时间: 2025-12 · arXiv:2512.18194 (preprint)

原文: https://arxiv.org/abs/2512.18194

作者/机构: Dongha Yoon (Virginia Tech), Younghoon Min (SK Hynix America), Hoshik Kim (SK Hynix America), Sam H. Noh (Virginia Tech), Jongryool Kim (SK Hynix America)

速读

一句话结论
TraCT 提出了一个基于 CXL 共享内存的机架级大语言模型推理系统,通过让 GPU 直接跨节点读写 KV Cache,彻底消除了分离式推理中的网络传输开销,在真实负载下将首字延迟降低了最多 9.8 倍,吞吐量提升了 1.6 倍。

要解决什么问题
当前的大语言模型服务通常采用分离式架构,将计算密集的 Prefill(预填充)阶段和对延迟敏感的 Decode(解码)阶段分配给不同的工作节点,以此提升资源利用率。这种架构引入了一个核心卡点:Prefill 节点生成的 KV 张量必须传输给 Decode 节点。现有的系统均依赖 RDMA 等网络协议栈进行跨节点交换。这意味着每一次 KV 数据的流转,都需要经过发送端主机内存、网卡队列、网络传输、接收端网卡和接收端主机内存。即使引入了全局前缀缓存(如 LMCache 或 Mooncake),缓存命中的 KV 块依然逃不掉这趟网络之旅。随着模型参数和上下文长度的增加,每次请求动辄产生数百兆字节的 KV 数据,这种跨网卡的序列化开销、额外的内存拷贝以及网络拥塞,不仅严重拖慢了首字延迟,导致长尾延迟极不稳定,也成为了限制整个系统吞吐量上限的物理瓶颈。

怎么做的
TraCT 的核心思路是将 CXL(Compute Express Link)Type-3 设备的共享内存池同时作为无网络的 KV 传输介质和机架级的全局前缀缓存。Prefill 节点生成未命中的 KV 块后,直接通过 GPU-CXL DMA 写入共享内存;Decode 节点则通过 DMA 直接从共享内存拉取到本地 GPU,全程绕过 CPU 网络栈和网卡。为了实现真正的零拷贝,TraCT 利用 CUDA 的主机内存注册接口将整个 CXL 内存区域锁定,防止驱动程序在主机 DRAM 中分配中间反弹缓冲区。为了在缺乏跨节点硬件原子操作和全局缓存一致性的 CXL 硬件上实现这一机制,TraCT 设计了三个关键部件。第一是双层软件同步锁:节点内进程先竞争本地 DRAM 锁,获胜者再去轮询 CXL 内存中的全局锁槽位,由一个后台锁管理器统一仲裁,从而将跨节点竞争的参与者数量严格限制在节点数以内。第二是基于偏移量的共享内存分配器:由于不同节点的虚拟地址基址不同,TraCT 摒弃了绝对指针,所有共享数据结构均采用偏移量寻址:

$$ \mathtt{ptr} = \mathtt{base} + \mathtt{off} $$


第三是非一致性内存下的可见性管理:由于 GPU-CXL DMA 直接绕过 CPU 缓存,大容量的 KV 负载本身不需要刷新;但对于前缀树索引等元数据,TraCT 采用细粒度的 clflush 指令强制将修改后的缓存行刷入 CXL 设备。为了避免复杂的树形结构引发高频的元数据更新和锁竞争,TraCT 采用固定大小的哈希表结合线性探查来管理前缀缓存,利用当前块的 Token 集合 $T_i$ 和前驱块哈希 $h_{i-1}$ 迭代计算块哈希:

$$ h_i = \mathrm{hash}(h_{i-1}, T_i) $$
只有在 DMA 传输彻底完成后,系统才会更新哈希表元数据并执行 clflush,以此作为数据全局可见的屏障。

效果如何
实验基于两台服务器搭建,配备 NVIDIA A6000 GPU、100Gbps 网卡和 64GB 的 Niagara 2.0 CXL Type-3 内存扩展器,运行 DeepSeek-R1-Distill-Llama-8B 模型。对比基线有两个:NIXL/UCX(代表无缓存的纯 RDMA 网络传输路线)和 LMCache(代表基于主机 DRAM 和网络传输的全局 KV 缓存路线)。在输入长度达 6000 Token 的静态负载下,TraCT 证明了纯 CXL 传输的延迟远低于 RDMA。在基于真实请求分布的合成负载测试中,TraCT 在 3.0 QPS 的压力下,相比 LMCache 将峰值吞吐量提升了 1.6 倍;平均首字延迟降低了最多 9.8 倍,P99 长尾延迟降低了 6.2 倍,同时显著降低了 GPU 的功耗和显存占用时间。该方法的代价在于,为了保证元数据更新的绝对正确性,必须使用同步的 clflush 而非开销更小的异步 clflushopt 指令;同时,为了避免跨 CPU 插槽的延迟惩罚,系统必须将所有工作线程严格绑定在直连 CXL 设备的 NUMA 节点上;此外,受限于元数据更新的同步开销,系统目前只能采用最简单的 LRU 淘汰策略,无法轻易引入更复杂的缓存替换算法。

主要贡献

现代大型语言模型(LLMs)服务系统正越来越多地采用分离式架构,将计算密集的Prefill(预填充)阶段与对延迟敏感的Decode(解码)阶段分离开来,以此提高资源利用率和吞吐量。然而,这种架构引入了一个核心瓶颈:在Prefill阶段生成的Key/Value (KV) 张量必须传输给Decode worker。现有的系统绝大多数依赖基于RDMA的网络路径来进行这种数据交换。随着模型规模和上下文长度的增加,每次请求交换的KV数据量高达数百兆字节,使得KV传输成为影响请求延迟和系统峰值吞吐量的主要因素。即使在使用前缀缓存复用技术时,现有的系统依然需要通过网卡队列、主机DRAM缓冲区以及复杂的传输层协议来搬运KV块,这种网络跳数(NIC hop)显著增加了Prefill延迟并限制了整体性能。

为了解决这一核心问题,本文提出了TraCT,这是一个机架级(rack-scale)的LLM服务系统。TraCT的研究目标是利用Compute Express Link (CXL) 共享内存既作为无网络的KV传输介质,又作为机架级的前缀感知KV缓存。其核心创新点包括:

  1. 网络无感知的KV传输与复用:TraCT允许GPU通过CXL的load/store和直接内存访问(DMA)操作,直接读写KV块,彻底消除了限制现有分离式流水线的网卡跳数。
  2. 针对非一致性CXL内存的软件同步机制:当前CXL设备不提供跨节点的硬件原子操作,也不保证全设备的硬件缓存一致性。TraCT创新性地提出了一种两层(two-tier)节点间软件锁同步机制,用以在无硬件原子操作的情况下保证互斥。
  3. 细粒度可见性与共享对象管理:为了在非一致性共享内存上保证元数据和数据的正确可见性,TraCT引入了细粒度的缓存行刷新策略,并设计了专为跨节点共享定制的内存分配器和基于偏移量的紧凑对象存储。

背景知识与设计原则

CXL共享内存基础与挑战。Compute Express Link (CXL) 【1,CXL specification+URL: https://computeexpresslink.org/cxl-specification/ 】 建立在PCIe之上,在CPU和设备之间提供高带宽、低延迟的通信。特别是CXL Type-3设备,通过CXL.mem协议提供主机可访问的内存。在当前系统中,CXL Type-3设备通常作为一个支持DAX的内存区域暴露出来,可以映射到用户空间作为字节寻址内存。从单节点软件的角度来看,它的行为类似于本地DRAM,支持正常的load/store和DMA读写。然而,当多个主机连接到同一个CXL设备时,硬件并不会自动在主机间保持缓存一致性。近期的研究表明,即使存在一致性,也仅限于特定设备的小区域,无法扩展到设备的全部容量 【8,Tigon: a distributed database for a cxl pod+2025+OSDI】,【9,Memory sharing with cxl: Hardware and software design approaches+2024】。因此,将CXL用作共享内存池的应用程序必须显式地管理跨节点的同步和数据可见性 【8,Tigon: a distributed database for a cxl pod+2025+OSDI】,【21,Partial failure resilient memory management system for (cxl-based) distributed shared memory+2023+SOSP】。CXL为机架级系统提供了无网络栈的内存通信语义,但缺乏跨节点原子操作和全设备一致性,这要求软件必须仔细处理缓存行状态、刷新和排序。

LLM推理的两阶段特性与KV缓存。现代LLM(如GPT、Gemini和Llama 【7,Gemini 2.5: Pushing the frontier with advanced reasoning, multimodality, long context, and next generation agentic capabilities+2025】,【11,The llama 3 herd of models+2024】,【12,Chatgpt (gpt-5)+URL: https://openai.com/chatgpt】)是仅解码器(decoder-only)的Transformer,自回归地生成token。推理分为Prefill和Decode两个阶段。在Prefill期间,长度为$N$的完整输入提示在一次前向传播中被处理,计算$O(N^2)$的注意力交互,此阶段是计算密集型的。模型生成第一个输出token以及$N$个K/V张量,形成初始KV缓存。在Decode阶段,模型逐个生成token,在步骤$t$时,模型仅计算新token的Q/K/V向量,并与之前存储在KV缓存中的$N+t-1$个张量进行注意力计算,此阶段转变为内存受限。为了避免重计算,系统缓存每一层的中间K和V张量。对于大模型,KV缓存占据了极大的内存(例 如Llama3 405B每token需$504\mathrm{KB}$)。

LLM推理中KV张量的演变
LLM推理中KV张量的演变

KV管理与分离式架构的耦合。在单节点内,系统如vLLM通过操作系统分页机制解决GPU内存碎片化,SGLang引入RadixAttention实现前缀共享 【10,Efficient memory management for large language model serving with pagedattention+2023+SOSP】,【22,Sglang: efficient execution of structured language model programs+2025+NIPS】。在多节点分离式架构中(如Splitwise、DistServe、Preble和Dynamo 【3,NVIDIA Dynamo+URL: https://github.com/ai-dynamo/dynamo】,【13 ,Splitwise: Efficient generative llm inference using phase splitting+2024+ISCA】,【17,Preble: Efficient distributed prompt scheduling for LLM serving+2025+ICLR】,【23,Distserve: disaggregating prefill and decoding for goodputoptimized large language model serving+2024+OSDI】),Prefill和Decode被拆分到不同worker,这改善了资源扩展,但每次请求都依赖RDMA等网络路径将KV块从Prefill传输到Decode worker。即使LMCache 【6,Lmcache: An efficient kv cache layer for enterprise-scale llm inference+2025】 和Mooncake 【14,Mooncake: Trading more storage for less computation — a KVCache-centric architecture for serving LLM chatbot+2025+FAST】 支持跨节点KV复用,它们依然保留了基于网络的KV张量移动。TraCT探索使用CXL共享内存直接作为机架级KV缓存和传输层,消除网络跳数。

TraCT的设计目标与挑战。TraCT旨在实现三个目标:1. 无网络的KV传输,用GPU与CXL间的直接DMA替代RDMA;2. 机架级KV复用,将KV块及前缀缓存元数据直接存入CXL共享内存;3. 去中心化KV管理,避免集中式元数据服务器,worker通过load/store直接管理状态。实现这些目标面临三大挑战:首先,必须在没有硬件原子操作和全局一致性的情况下确保共享内存区域的互斥访问。现有方法要么依赖极小的硬件一致性区域,要么使用消耗$O(N^2)$内存的队列(如cMPI 【18,cmpi: Using cxl memory sharing for mpi one-sided and two-sided inter-node communications+2025+SC】),或者引入违背CXL初衷的集中式服务器(如Beluga 【19,Beluga: A cxlbased memory architecture for scalable and efficient llm kvcache management+2025】)。其次,在缺乏硬件缓存一致性的情况下,必须处理缓存刷新以保证数据可见性。最后,需要一种不依赖进程私有虚拟地址指针的共享对象管理机制。

方法细节

TraCT架构概览。在TraCT架构中,每个机架包含多个Prefill和Decoding worker,它们均连接到一个共享的CXL Type-3设备。TraCT将该CXL设备暴露为所有参与服务器上支持DAX映射的字节寻址区域。TraCT底层库提供三个原语:节点间锁、内存分配器和对象存储。基于该库,TraCT实现了前缀缓存索引和GPU-CXL KV块传输模块。当LLM请求到达时,Prefill worker首先查找前缀缓存,若命中则执行CXL到GPU的KV块传输;随后计算缺失的KV块,将前缀缓存条目发布并将缺失的KV块通过GPU-to-CXL DMA写入CXL,最后释放GPU资源。对于Decoding worker,在接收到请求后,它直接从CXL读取该提示的所有KV块到GPU内存中,随后开始逐token的生成计算。

TraCT概览
TraCT概览

确保互斥访问的两层锁定结构。由于CXL Type-3设备不支持跨节点原子指令,TraCT无法依赖传统的硬件锁。如果让所有节点的所有进程直接竞争一个全局锁,由于进程动态加入和退出会导致参与者数量无界,产生极高的竞争。因此,TraCT采用了一种两层(two-tier)软件同步结构,由成对的global_lock和local_lock组成。

TraCT中的两层节点间锁定
TraCT中的两层节点间锁定

保证数据可见性的缓存一致性策略。由于没有跨主机的硬件一致性,节点可能会读取到私有缓存中的过期数据。TraCT通过以下策略保证正确性:

TraCT软件栈
TraCT软件栈

共享内存分配器与对象存储的实现。为了管理CXL中的字节寻址区域,TraCT设计了内存分配器和对象存储:

前缀缓存管理的静态设计。动态的树结构在插入或分裂时会引发频繁的指针更新和缓存行刷新,在非一致性CXL内存中开销极高。TraCT采用固定大小的哈希表结合线性探测来管理前缀缓存。

共享内存中的数据结构实现细节。由于不同节点操作系统映射CXL区域的虚拟基址不同,TraCT在所有共享数据结构中采用基于偏移量的寻址方式。每个节点保存其本地虚拟基址,进行如下转换:$\mathtt{ptr} = \mathtt{base} + \mathtt{off}$,$\mathtt{off} = \mathtt{ptr} - \mathtt{base}$。同时,共享对象严格按缓存行边界对齐,以避免伪共享。

启用直接GPU-CXL DMA与NUMA优化。直接调用cudaMemcpy()会导致CUDA驱动在主机DRAM中分配中间反弹缓冲区。TraCT利用CUDA的主机内存注册接口,将整个CXL共享内存区域固定(pin)。这使得CUDA运行时将其视为页锁定主机内存,允许DMA引擎直接访问CXL设备,实现真正的零拷贝。此外,为了避免跨CPU插槽的流量,TraCT将所有相关线程(包括锁管理器和KV连接器线程)绑定到直接连接CXL设备的NUMA节点上。

实验环境

实验结果

实验1:CXL作为传输介质的有效性(无缓存)

TraCT与NIXL TTFT CDF对比
6000 token输入下的吞吐量对比

实验2:端到端吞吐量与延迟性能

峰值吞吐量
工作负载的前缀缓存命中率
TTFT CDF

实验3:单请求时间分解与GPU资源分析

每个请求的时间分解
QPS=3.0下GPU资源消耗随时间变化

结论

本文提出了TraCT,这是一个机架级的LLM服务系统,它使用直接的GPU-CXL DMA取代了基于RDMA的KV传输,使得Prefill和Decode worker能够通过CXL共享内存高效地共享KV块。为了在非一致性的CXL Type-3设备上正确运行,TraCT引入了两层节点间锁、软件管理的数据可见性机制,以及专为节点间共享设计的分配器和对象存储。在真实的CXL硬件上基于Dynamo-vLLM运行时的评估表明,TraCT相比于基于网络的基线系统,显著改善了TTFT和峰值吞吐量,证明了CXL共享内存是分离式LLM服务的一个实用且高性能的基础底座。

补充细节

相关工作对比分析。


方法细节中的引用汇总: