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

arXiv:2512.18194 arXiv 2025 CXL Shared Memory Disaggregated Serving 9.8x TTFT Reduction 6.2x P99 Reduction

Dongha Yoon, Younghoon Min, Hoshik Kim, Sam H. Noh, Jongryool Kim

Virginia Tech, SK Hynix America

Overview

Figure 1: Evolution of KV tensors in LLM inference
Figure 1 (from paper): Evolution of KV tensors in LLM inference. Shows how KV cache management has evolved from monolithic to disaggregated architectures, motivating the need for efficient cross-node KV transfer.

The Network Bottleneck in Disaggregated Serving

Disaggregated LLM serving (DistServe, Splitwise, Dynamo) separates prefill and decode into distinct workers. The KV tensors generated during prefill must be transferred to decode workers. Current systems use RDMA over network, which traverses NIC queues, host DRAM buffers, and layered transport protocols. This network hop inflates prefill latency, increases tail variability, and constrains throughput—even with prefix reuse, every cache hit still requires network transfer.

TraCT: Replace RDMA with CXL Shared Memory

TraCT uses CXL Type-3 shared memory as both a network-free KV transfer substrate and a rack-wide prefix-aware KV cache. GPUs write KV blocks directly into CXL memory via GPU-CXL DMA and fetch them back the same way—no NIC involvement, no host-to-host copies, no network serialization. This eliminates the RDMA hop entirely, turning prefill-decode communication into local load/store and DMA operations on a shared address space.

Architecture

Figure 2: TraCT overview
Figure 2 (from paper): TraCT overview. Prefill and decode workers communicate KV tensors through CXL shared memory via direct GPU-CXL DMA, eliminating the RDMA network hop entirely.
TraCT System Overview
Prefill Worker
GPU + NVIDIA A6000
KV Write (GPU → CXL DMA)
CXL Shared Memory
Prefix Cache
+ KV Blocks
Niagara 2.0 CXL Type-3
Decode Worker
GPU + NVIDIA A6000
KV Read (CXL → GPU DMA)
No NIC, no RDMA, no network stack. Direct GPU-CXL DMA over PCIe/CXL fabric.

Key Mechanisms

Figure 3: Two-tier inter-node locking in TraCT
Figure 3 (from paper): Two-tier inter-node locking in TraCT. Processes on each node first contend on its local DRAM-resident lock, so that at most one thread per node participates in acquiring the global lock. The global lock resides in CXL memory with per-node slots in three states: I (idle), W (waiting), L (lock granted).

Two-Tier Inter-Node Locking

CXL Type-3 devices lack cross-node atomic instructions. TraCT implements a two-tier software lock: local_lock (per-node DRAM, pthread_mutex) ensures at most one process per node contends globally, and global_lock (CXL shared memory) coordinates across nodes. A lock manager thread polls global_lock entries, bounding contention to the number of nodes (typically 2-4 in a rack).

Cache Coherence via clflush

CXL shared memory is non-coherent across nodes. TraCT uses clflush (not clflushopt) to ensure metadata visibility: it evicts cachelines from the local hierarchy before the instruction completes, guaranteeing that other nodes observe the update. KV payloads bypass CPU caches entirely via GPU-CXL DMA, requiring no software coherence.

Shared Memory Allocator & Object Store

Two-tier allocator: global chunk allocator (CXL bitmap) + per-node heap allocators (local DRAM free-lists). Shared object store publishes only root objects (prefix index hash table) via offsets, not pointers. Offset-based addressing ensures consistency across nodes with different virtual address mappings.

Direct GPU-CXL DMA

CXL shared memory is registered via CUDA's host-memory registration interface and pinned as page-locked host memory. This enables true zero-copy GPU-CXL transfers via cudaMemcpy without intermediate host DRAM bounce buffers. Cache-hit KV blocks are consumed directly from CXL with no regeneration or network transfer.

// TraCT KV transfer: prefill writes, decode reads // Prefill worker: kv_block = gpu_compute_kv(prompt, layer) cxl_addr = shmalloc(kv_block.size) # CXL allocator cudaMemcpy(cxl_addr, gpu_kv_ptr, size, DeviceToHost) # GPU->CXL DMA cxl_shm_acquire_lock(prefix_lock) publish_prefix_entry(block_hash, cxl_addr) clflush(prefix_entry) # ensure visibility cxl_shm_release_lock(prefix_lock) // Decode worker: cxl_addr = lookup_prefix(block_hash) # hash table lookup cudaMemcpy(gpu_kv_ptr, cxl_addr, size, HostToDevice) # CXL->GPU DMA

Performance Results

9.83x
Average TTFT Reduction
6.2x
P99 TTFT Reduction
1.6x
Peak Throughput Improvement
~5K
Lines of C/C++
SystemAvg TTFTP99 TTFTPeak ThroughputCache Backend
NIXL/UCX (RDMA)1.0x (baseline)1.0x1.0xNone (transfer only)
LMCache (DRAM cache)~0.3x~0.4x~1.0xHost DRAM (48GB)
TraCT (no cache)~0.7x~0.8x1.6xCXL (transfer only)
TraCT (with cache)0.10x (9.83x better)0.16x (6.2x better)1.6xCXL shared memory (64GB)
Hardware: Two servers with NVIDIA A6000 GPUs (48GB GDDR), 512GB host DRAM, connected via Niagara 2.0 CXL Type-3 memory expander (64GB, 640ns access latency, 10.1 GB/s bandwidth). Model: DeepSeek-R1-Distill-Llama-8B. Even without caching, CXL-based DMA provides lower and more stable TTFT than RDMA due to eliminated NIC queuing and transport variability.
Figure 5: TraCT vs NIXL TTFT CDF
Figure 5 (from paper): TraCT vs NIXL TTFT CDF. The number indicates input token length. TraCT achieves significantly lower and tighter TTFT distributions across all input sizes.
Figure 6: TraCT (no cache) vs NIXL throughput
Figure 6 (from paper): TraCT (no cache) vs. NIXL throughput for 6000-token inputs. Even without prefix caching, CXL-based transfer delivers higher throughput than RDMA.
Figure 7: Peak throughput
Figure 7 (from paper): Peak throughput comparison across systems. TraCT achieves 1.6x higher peak throughput than RDMA-based baselines.
Figure 8: Prefix cache hit rate of workloads
Figure 8 (from paper): Prefix cache hit rate of workloads. Higher hit rates translate directly to greater TTFT reduction with TraCT's CXL shared cache.
Figure 9: TTFT CDF
Figure 9 (from paper): TTFT CDF showing the distribution of time-to-first-token across different system configurations with prefix caching enabled.
Figure 10: Per-request time breakdown
Figure 10 (from paper): Per-request time breakdown showing how TraCT redistributes latency components compared to RDMA-based systems.
Figure 11: GPU resource consumption over time
Figure 11 (from paper): GPU resource consumption over time for prefill and decode workers under QPS = 3.0. RX denotes GPU receive bandwidth (PCIe/CXL to GPU), and TX denotes GPU transmit bandwidth (GPU to PCIe/CXL).

Workload Sensitivity Analysis

Prefix Reuse Rate

TraCT's gains are most pronounced when the prefix reuse rate is high. Workloads like multi-turn chat (identical system prompt across turns), RAG (identical document context for multiple user queries), or agentic pipelines (identical tool definitions across calls) benefit most. For workloads with fully unique prompts (e.g., code generation with entirely new inputs each time), the TraCT no-cache configuration applies — which still provides ~3x lower transfer latency than RDMA due to eliminated NIC overhead.

Sequence Length Impact

Longer input sequences generate more KV blocks, increasing both the benefit and the cost of caching. For short inputs (512 tokens), the absolute KV data size is small and RDMA latency is dominated by fixed overhead — CXL's advantage is primarily in eliminating that fixed NIC overhead. For long inputs (4K–32K tokens), the total KV size grows (hundreds of MB for 70B models), making the 10.1 GB/s CXL bandwidth the primary constraint. The sweet spot for current TraCT hardware is 1K–8K token sequences with 7B–13B models.

Significance

First CXL-Based LLM Serving System

TraCT is the first system to demonstrate rack-scale KV caching and transfer for disaggregated LLM serving on real CXL Type-3 hardware (Niagara 2.0). It proves CXL shared memory is a practical high-performance substrate for LLM inference.

Eliminates the Network Hop

By replacing RDMA with direct GPU-CXL DMA, TraCT removes the NIC, host DRAM buffers, and transport protocol overhead from the KV transfer critical path. This produces tighter, more predictable TTFT distributions.

Software Solutions for Non-Coherent Memory

Addresses three fundamental challenges of multi-host CXL (no atomic operations, no coherence, no shared pointers) with practical software mechanisms: two-tier locking, clflush-based visibility, and offset-based addressing.

Energy and GPU Efficiency

TraCT lowers GPU SM utilization and PCIe bandwidth usage during both prefill and decode. Cache hits bypass KV regeneration entirely. These translate to lower power consumption and reduced total cost of ownership.

CXL Memory Architecture Deep Dive

What is CXL 3.0 Shared Memory?

Compute Express Link (CXL) is a high-speed, low-latency interconnect built on top of the PCIe 5.0 physical layer. CXL 3.0 introduces Fabric Manager support and multi-host shared memory semantics. A CXL Type-3 memory device (e.g., Niagara 2.0) exposes a large DRAM pool that can be mapped into the physical address space of multiple host servers simultaneously. Unlike RDMA, which sends data through a NIC, operating system network stack, and transport queues, CXL memory is accessed via ordinary CPU load/store instructions and GPU DMA — making it feel like local DRAM to the processor.

Three-Tier Memory Hierarchy in TraCT

💻
Tier 1 — GPU VRAM (On-Device)
Capacity: 48 GB per GPU (NVIDIA A6000) Latency: ~100 ns BW: ~768 GB/s Role: Active KV blocks for current batch
🧠
Tier 2 — CXL Shared Memory Pool (TraCT Target)
Capacity: 64 GB (Niagara 2.0, shared across rack) Latency: ~640 ns BW: ~10.1 GB/s per node Role: Shared KV cache + in-flight KV transfer
💾
Tier 3 — Host CPU DRAM (Overflow / Metadata)
Capacity: 512 GB per node Latency: ~80 ns (local) Role: Local locks, per-node allocator free-lists, non-shared control plane

CXL vs RDMA: Path Comparison

RDMA KV Transfer Path (traditional disaggregated serving):

GPU VRAM
Host DRAM buf
NIC + Net Stack
Remote DRAM
Decode GPU

TraCT CXL Transfer Path:

Prefill GPU VRAM
CXL Shared Memory
Decode GPU VRAM

3 hops eliminated. No NIC, no OS network stack, no transport protocol.

Why Not Just Use Faster Network?

Even with InfiniBand HDR (200 Gb/s), the KV transfer bottleneck is not raw bandwidth — it is latency and queuing variability. For a single KV block (e.g., 4 KB for a 32-layer model with GQA), PCIe transfer takes tens of microseconds. CXL latency is ~640 ns, roughly 30–100x lower depending on block size and NIC queue depth. Moreover, NIC queues serialize requests from multiple concurrent LLM requests, creating head-of-line blocking that inflates P99 TTFT. CXL memory has no such serialization point.

Latency Breakdown: Transfer Methods

Relative transfer latency (lower = better; RDMA baseline = 1.0)
RDMA (InfiniBand)
1.0x
baseline
KV Recompute (no cache)
~0.75x
similar
CXL DMA (no cache)
~0.3x
3x faster
CXL DMA (cache hit)
0.10x
9.8x faster

Cache-hit path: KV is already in CXL pool from a previous request with the same prefix. DMA to decode GPU VRAM bypasses all compute.

Request Lifecycle in TraCT

Cache-Miss Path (first request)

1. Request arrives at Prefill Worker
Scheduler dispatches prompt tokens
2. GPU computes KV tensors (all layers)
Standard transformer prefill, O(n²) attention
3. cudaMemcpy: GPU VRAM → CXL pool
GPU-CXL DMA, ~10.1 GB/s, no CPU involvement
4. Publish prefix entry in CXL hash table
Acquire global_lock → write offset → clflush → release
5. Decode Worker: lookup → CXL → GPU DMA
cudaMemcpy CXL → GPU, then autoregressive decode begins

Cache-Hit Path (repeated prefix)

1. Request with matching prefix arrives
System prompt, few-shot examples, tool definitions reused
2. Scheduler finds hit in CXL prefix index
Hash lookup on block_hash over CXL memory-mapped table
3. Prefill Worker: compute only NEW tokens
Prefix KV already in CXL — GPU skips recompute for those layers
4. Decode Worker fetches KV directly from CXL
GPU-CXL DMA without any network transfer
5. Decode begins — 9.8x faster TTFT
No NIC traversal, no OS stack, no queuing jitter
Key design decision: TraCT's prefix cache uses block-level granularity (same as PagedAttention). Each block is independently addressable in CXL via an offset in the shared hash table. This means partial prefix matches are still useful — if 80% of a system prompt prefix is already cached, only the remaining 20% needs recomputation.

Memory Management: Eviction Policy

CXL Pool Eviction

When the 64 GB CXL pool is full, TraCT evicts blocks using an LRU (Least Recently Used) policy at the block level. The eviction is coordinated through the two-tier locking protocol — the global_lock must be acquired to remove a prefix entry from the shared hash table. Evicted blocks are not spilled to CPU DRAM; they are simply freed. A future request for that prefix will trigger a cache miss and full prefill recomputation.

GPU VRAM Pressure Relief

A secondary benefit of CXL offloading is GPU VRAM pressure relief. In a non-disaggregated serving system, the KV cache must share GPU VRAM with model weights and activations. With TraCT, KV blocks for completed or long-idle requests can be offloaded from GPU VRAM to CXL, freeing VRAM for active batches. This is similar to how vLLM uses CPU DRAM swapping, but with 4–8x lower latency than DRAM DMA over PCIe.

Hardware Architecture & Rack-Scale Deployment

Experimental Hardware Setup

Component Specification Role in TraCT
GPU NVIDIA RTX A6000 (48 GB GDDR6) Prefill compute + decode KV attention
Host DRAM 512 GB DDR5 per node Local locks, allocator free-lists, process memory
CXL Memory Device Niagara 2.0 (CXL Type-3, 64 GB) Shared KV cache pool + in-flight KV staging
CXL Latency ~640 ns (round-trip read) vs ~80 ns local DRAM; vs ms-level RDMA
CXL Bandwidth 10.1 GB/s per node Sufficient for most KV block sizes; shared across nodes
LLM Model DeepSeek-R1-Distill-Llama-8B Evaluation model; fits on single A6000 per worker
Server Count 2 nodes (1 prefill + 1 decode) Minimum rack-scale CXL topology

Rack-Scale CXL Topology

CXL 3.0 Shared Memory Fabric — Rack View
Prefill Node 1
Server A
NVIDIA A6000
48 GB GDDR6
512 GB DDR5
Prefill Node 2
Server B
NVIDIA A6000
48 GB GDDR6
512 GB DDR5
Decode Node 1
Server C
NVIDIA A6000
48 GB GDDR6
512 GB DDR5
Decode Node 2
Server D
NVIDIA A6000
48 GB GDDR6
512 GB DDR5
↓ ↓ ↓ ↓
CXL 3.0 Switch / Fabric Manager
↓
CXL Type-3 Memory Device
Niagara 2.0
64 GB Shared DRAM Pool
~640 ns latency • 10.1 GB/s per node
Shared by all nodes in rack
All nodes share the same CXL physical address range. Each node maps the CXL pool into its virtual address space. No NIC or network hop between nodes for KV data.

CXL vs PCIe: Connectivity Model

CXL is built on the PCIe 5.0 physical layer but extends it with three protocols: CXL.io (device discovery, config), CXL.cache (device-side cache coherency with host), and CXL.mem (host-managed device memory). The key difference from plain PCIe: CXL.mem exposes device DRAM as a coherent memory extension of the host's address space, allowing load/store semantics rather than explicit DMA commands for metadata. For KV payload, TraCT still uses explicit GPU DMA for throughput.

Scaling Considerations

The paper demonstrates a 2-node setup (1 prefill + 1 decode). CXL 3.0 fabric can support up to 4,096 nodes in theory, though practical rack-scale deployments target 16–32 nodes. The two-tier locking overhead grows with node count: locking contention is bounded by the number of nodes, not concurrent requests. For 4 nodes, lock overhead is negligible; for 32 nodes, careful tuning of lock manager polling intervals may be needed.

Hardware Availability Note: CXL Type-3 memory devices capable of multi-host shared memory (like Niagara 2.0) are currently available from SK Hynix, Samsung, and Micron, but require server motherboards with CXL 2.0+ support and a CXL switch for multi-host topologies. As of 2025, CXL switches from Microchip (Switchtec PFX) and Astera Labs are commercially available.

Limitations & Future Work

Bandwidth Ceiling

The Niagara 2.0 device provides ~10.1 GB/s per node. For large models (e.g., Llama-3 70B with 80 layers), a single full-sequence KV cache can exceed 1 GB. Under high request rates, the shared 10.1 GB/s becomes a bottleneck. The paper evaluates on a smaller 8B model; scaling to 70B+ models would require either multiple CXL devices, a CXL switch providing higher aggregate bandwidth, or selective caching of only the most frequently accessed layers.

No Hardware Atomics on CXL Type-3

CXL Type-3 devices do not support cross-host atomic operations. TraCT's two-tier software locking adds overhead for every prefix index write. For read-heavy cache-hit workloads, read operations are lock-free (only a hash table lookup), but any metadata write — including eviction — must go through the locking protocol. Future CXL revisions or CXL.cache semantics could enable hardware-assisted atomic metadata updates.

Evaluation on 2-Node Setup Only

The experimental evaluation uses a single prefill node and a single decode node. Real production disaggregated serving (e.g., Dynamo) runs with hundreds of nodes. The paper does not evaluate or characterize multi-node CXL contention, lock manager scalability, or how the CXL bandwidth degrades as more nodes simultaneously access the shared pool. These are open questions for follow-up work.

Integration with Existing Serving Frameworks

TraCT is implemented as a standalone prototype (~5K lines C/C++) rather than integrated into vLLM, TensorRT-LLM, or SGLang. The paper demonstrates feasibility but does not provide a production-ready integration path. Future work could implement TraCT as a vLLM KV cache connector plugin, analogous to how MooncakeTransferEngine or LMCache integrate. The CXL shared memory layer would sit below the PagedAttention block manager.

Future Directions

Near-term (hardware-side):

  • CXL 3.1 hardware atomics support
  • Higher-bandwidth CXL devices (>50 GB/s)
  • NVLink-CXL bridge for multi-GPU nodes
  • Larger CXL pools (multi-TB at rack scale)

Near-term (software-side):

  • vLLM/SGLang plugin integration
  • Semantic/semantic-aware prefix hashing
  • Speculative eviction policies (predict future reuse)
  • Fault tolerance: CXL failure fallback to RDMA
Contextual positioning: TraCT is most relevant to hyperscale data centers that are building CXL-equipped racks for disaggregated memory infrastructure. The system is complementary to, not a replacement for, RDMA-based KV transfer in multi-rack or cross-datacenter deployments — CXL is inherently rack-scoped. In environments where all prefill/decode workers fit within a single rack, TraCT provides a compelling improvement with no network stack changes required.

Implementation Deep Dive

Figure 4: TraCT software stack
Figure 4 (from paper): TraCT software stack. The userspace library provides CXL shared memory allocation, prefix index management, and GPU-CXL DMA wrappers, integrated with the vLLM serving framework.

TraCT Software Stack (~5,000 lines C/C++)

TraCT is implemented as a userspace library that prefill and decode workers link against. The library provides three primary interfaces: the CXL shared memory allocator, the prefix index (hash table over CXL), and the GPU-CXL DMA wrappers. The ~5K line count excludes the vLLM Python integration shim.

/* TraCT public API sketch (simplified) */ /* 1. Initialize: map CXL device into process VA space */ tract_ctx_t* tract_init( const char* cxl_dev_path, /* e.g., /dev/dax0.0 */ size_t pool_size, /* 64 GB */ int node_id /* 0 = prefill, 1 = decode */ ); /* 2. Allocate a KV block in CXL pool */ cxl_offset_t tract_alloc(tract_ctx_t* ctx, size_t size); /* 3. GPU → CXL DMA (prefill worker writes KV) */ cudaError_t tract_gpu_to_cxl( tract_ctx_t* ctx, cxl_offset_t dst, void* gpu_src, size_t size ); /* 4. Publish prefix to shared index */ void tract_publish_prefix( tract_ctx_t* ctx, uint64_t block_hash, /* hash of token sequence */ cxl_offset_t kv_offset ); /* 5. Lookup prefix in shared index (lock-free read) */ cxl_offset_t tract_lookup_prefix(tract_ctx_t* ctx, uint64_t block_hash); /* 6. CXL → GPU DMA (decode worker loads KV) */ cudaError_t tract_cxl_to_gpu( tract_ctx_t* ctx, void* gpu_dst, cxl_offset_t src, size_t size );

CUDA Host Memory Registration

For GPU-CXL DMA to work without bounce buffers, the CXL memory region must be registered with CUDA as pinned (page-locked) host memory. TraCT calls cudaHostRegister(cxl_base_ptr, pool_size, cudaHostRegisterDefault) at initialization. This tells the CUDA runtime that the memory is host-accessible and that the GPU can DMA to/from it directly via PCIe without staging through GPU-attached DRAM. The registration is process-local: each node registers its own mapping of the CXL physical range.

Offset-Based Pointer Sharing

A fundamental challenge with shared memory across processes and machines is that virtual address mappings differ. Node A might map the CXL pool at virtual address 0x7f0000000000, while Node B maps it at 0x7e0000000000. Raw pointers stored in the shared hash table would be meaningless across nodes. TraCT solves this by storing CXL offsets (byte distance from the start of the pool) rather than virtual addresses. Each node translates: virtual_addr = cxl_base + offset. The base pointer differs per process, but the offset is universal.

Two-Tier Lock: Step-by-Step

When a prefill worker wants to publish a new prefix entry, it must acquire both locks in order:

/* Step 1: Acquire local_lock (node-local pthread_mutex in DRAM) */ pthread_mutex_lock(&node->local_lock); /* Step 2: Spin on global_lock in CXL shared memory */ while (!cxl_try_acquire(&cxl->global_lock[LOCK_PREFIX_IDX])) { cpu_relax(); /* busy-wait; bounded by node count */ } /* Step 3: Perform metadata update in CXL */ cxl->prefix_table[hash].offset = kv_offset; cxl->prefix_table[hash].valid = 1; /* Step 4: Force visibility to remote CPUs */ clflush(&cxl->prefix_table[hash]); /* Step 5: Release global_lock, then local_lock */ cxl_release(&cxl->global_lock[LOCK_PREFIX_IDX]); pthread_mutex_unlock(&node->local_lock); /* Reader side: lock-free (no lock needed for lookup) */ entry = cxl->prefix_table[hash]; if (entry.valid) prefetch(cxl_base + entry.offset); /* optional */

clflush (not clflushopt) is used because it serializes: the instruction does not complete until the cacheline is evicted from the local CPU cache hierarchy. clflushopt is asynchronous and would allow a remote node to observe a stale value if it reads immediately after.

Prefix Hashing: Block-Level Granularity

TraCT uses the same prefix hashing strategy as vLLM's RadixAttention and SGLang's prefix caching. The KV cache is divided into fixed-size blocks (e.g., 16 or 32 tokens per block). Each block is identified by a hash computed over all token IDs in the block concatenated with the hash of the preceding block. This chained hashing ensures that a block hash uniquely identifies both the content and position in the sequence. The CXL prefix index maps block_hash → CXL offset for every block currently resident in the pool.