Disaggregated Inference Architecture: Paged KV-Cache Tiering, Chunked Prefill, and Asynchronous Overlap
A deep-dive systems engineering guide on architecting disaggregated LLM serving pipelines using Paged KV-cache tiering, chunked prefill, and fused CUDA kernels.
Introduction
Modern large language model (LLM) serving infrastructure faces an inherent operational tension dictated by the asymmetric computational profiles of the autoregressive transformer. The execution lifecycle of a request decomposes into two distinct phases with orthogonal hardware bottlenecks:
- The Prefill Phase (Prompt Processing): Compute-bound and characterized by high arithmetic intensity ($\text{FLOPs} / \text{Byte}$). The model processes all input tokens concurrently, saturating tensor cores through large matrix multiplications (GEMMs).
- The Decode Phase (Token Generation): Memory-bandwidth-bound and characterized by low arithmetic intensity. The model generates tokens sequentially, issuing memory-bound matrix-vector (GEMV) operations that repeatedly fetch model weights and Key-Value (KV) cache tensors from High Bandwidth Memory (HBM) to on-chip SRAM.
In traditional monolithic serving engines (e.g., standard continuous batching systems), prefill and decode workloads are scheduled onto the same physical GPUs. This spatial and temporal colocation induces severe performance degradation:
- Resource Contention: Compute-intensive prefill bursts preempt or starve latency-critical decode steps, triggering high Inter-Token Latency (ITL) variance and tail latency amplification.
- Underutilization: Compute cores idle during decode-dominant batches due to memory bus saturation, while memory bandwidth remains underutilized during long-context prefills.
- Rigid Parallelism Topologies: Prefill scales efficiently with Context Parallelism (CP) and Tensor Parallelism (TP), whereas Decode scales better under Pipeline Parallelism (PP) and Data Parallelism (DP) with dynamic continuous batching.
Monolithic Serving (Interference & Resource Stranding)
+------------------------------------------------------------------+
| GPU HBM: [ Weights ] [ KV-Cache Batch A ] [ KV-Cache Batch B ] |
| Compute: [ Prefill GEMM (Compute Bound) ] <--- Contention ---> |
| [ Decode GEMV (Memory Bound) ] |
+------------------------------------------------------------------+
Disaggregated Serving (Phase-Specialized Compute Pools)
+-------------------------------+ +-------------------------------+
| Prefill Pool (P-Nodes) | | Decode Pool (D-Nodes) |
| - High Tensor Core FLOP/s | RDMA | - Massive Aggregate HBM BW |
| - Optimized for GEMM / TP / CP| ===> | - High Batch Sizes / DP |
| - Ephemeral KV Allocation | | - Tiered Paged KV-Cache |
+-------------------------------+ +-------------------------------+
Disaggregated LLM inference decouples prefill compute engines from decode serving pools across distinct physical hardware topologies connected via high-speed RDMA interconnects. This architectural paradigm optimizes hardware allocation for each phase, enforces strict Service Level Objectives (SLOs) on Time-to-First-Token (TTFT) and ITL independently, and unlocks global KV-cache lifecycle management.
Anatomy of Prefill-Decode Disaggregation
The mathematical rationale for disaggregation stems from the disparity in operational intensity between the two execution phases. Let $L$ be the sequence length, $d_{\text{model}}$ the hidden dimension, $n_{\text{layers}}$ the layer count, and $B$ the batch size.
$$\text{Arithmetic Intensity}{\text{prefill}} \approx \frac{24 \cdot B \cdot L \cdot n{\text{layers}} \cdot d_{\text{model}}^2}{2 \cdot n_{\text{layers}} \cdot d_{\text{model}}^2 + 4 \cdot B \cdot L \cdot n_{\text{layers}} \cdot d_{\text{model}}} \approx \mathcal{O}(L)$$
$$\text{Arithmetic Intensity}{\text{decode}} \approx \frac{24 \cdot B \cdot 1 \cdot n{\text{layers}} \cdot d_{\text{model}}^2}{2 \cdot n_{\text{layers}} \cdot d_{\text{model}}^2 + 4 \cdot B \cdot L \cdot n_{\text{layers}} \cdot d_{\text{model}}} \approx \mathcal{O}(1)$$
When $L \gg 1$, the prefill operational intensity exceeds the machine balance point of modern accelerators (e.g., $\approx 150 \text{ FLOP/Byte}$ on an NVIDIA H100 SXM5), fully saturating the 4th-generation Tensor Cores. Conversely, decode operational intensity hovers between $1$ and $10 \text{ FLOP/Byte}$, throttling hardware utilization to the memory bandwidth ceiling of the HBM3 subsystem ($3.35 \text{ TB/s}$).
Data Path and Transfer Topology
In a disaggregated architecture, when a request arrives at the ingress load balancer:
- Dispatch: The router directs the prompt context to a specialized Prefill Node (P-Node) configured with high tensor parallelism (e.g., $\text{TP}=8$ via NVLink).
- Execution & Context Generation: The P-Node processes the prompt, generates the initial token, and writes the intermediate Key and Value activation tensors directly into a contiguous staging buffer.
- KV-Cache Transport: The P-Node dispatches the generated KV-cache tensors over a non-blocking $800 \text{ Gbps}$ RoCEv2 or InfiniBand fabric to a target Decode Node (D-Node) using GPUDirect RDMA.
- Iterative Generation: The D-Node assigns the incoming KV tensor blocks into its local Paged KV-Cache allocator and executes autoregressive decode steps until sequence termination.
+---------------+ HTTP/gRPC +-----------------------+
| Client Ingress| ------------------> | Global Gateway/Router |
+---------------+ +-----------------------+
/ \
(1) Dispatch / \ (3) Assign Worker
Prompt v v
+--------------------+ +--------------------+
| Prefill Instance | | Decode Instance |
| (TP=8, NVLink) | | (TP=1/2, DP=N) |
+--------------------+ +--------------------+
| ^
| (2) GPUDirect RDMA Transfer |
| [KV Cache Tensor Blocks] |
+----------------------------------+
The primary engineering constraint of this architecture is the transport latency $T_{\text{transfer}}$ of the KV-cache relative to the model execution latency. For a request with prompt length $S$, batch size $1$, layer count $N_L$, number of KV heads $H_{KV}$, and head dimension $D$, the payload size in bytes $M_{\text{KV}}$ under FP8 precision is:
$$M_{\text{KV}} = 2 \times N_L \times H_{KV} \times S \times D \times 1 \text{ Byte}$$
For a 70B parameter model ($N_L = 80, H_{KV} = 8, D = 128$) with a $4096$-token prompt:
$$M_{\text{KV}} = 2 \times 80 \times 8 \times 4096 \times 128 \times 1 \approx 671.08 \text{ MB}$$
Over an $800 \text{ Gbps}$ ($100 \text{ GB/s}$ theoretical, $\approx 85 \text{ GB/s}$ effective) interconnect, the network serialization delay is:
$$T_{\text{transfer}} = \frac{671.08 \text{ MB}}{85 \text{ GB/s}} \approx 7.89 \text{ ms}$$
Because the prefill execution for $4096$ tokens on an 8-GPU H100 cluster requires approximately $35\text{--}45\text{ ms}$, network transfer overhead accounts for less than $20%$ of TTFT. This latency can be overlapped via pipelined chunk streaming.
Distributed Paged KV-Cache Tiering and Memory Layout
Memory fragmentation in naive transformer serving systems waste up to $60\text{--}80%$ of GPU memory through internal and external allocation gaps. Paged KV-cache abstractions solve this by allocating physical memory in fixed-size blocks (pages), virtualizing the address space of attention tensors.
In a disaggregated setup, this abstraction extends across a multi-tier storage hierarchy:
- Tier 1: GPU High Bandwidth Memory (HBM3e). Ultra-low latency ($\approx 100 \text{ ns}$ access latency, $>3\text{--}4.8 \text{ TB/s}$ per GPU), stores active decode sequences.
- Tier 2: Host Pinned System Memory (DRAM). Expanded capacity via PCIe Gen5 ($\approx 64 \text{ GB/s}$ bidirectional bandwidth per 16x slot), acts as a local spillover cache for paused or pre-allocated contexts.
- Tier 3: Remote Networked Storage Pool (CXL / NVMe-oF). Distributed block stores designed for prefix-cache persistence and cross-node migration.
Virtual Sequence Address Space (Request Alpha)
+-------------------+-------------------+-------------------+
| Logical Block 0 | Logical Block 1 | Logical Block 2 |
| (Tokens 0 - 15) | (Tokens 16 - 31) | (Tokens 32 - 47) |
+-------------------+-------------------+-------------------+
| | |
v v v
+-------------------+-------------------+-------------------+
| Physical Frame 42 | Physical Frame 09 | Physical Frame 88 |
| (GPU 0 HBM) | (GPU 0 HBM) | (Host RAM Tier 2) |
+-------------------+-------------------+-------------------+
Physical Memory Block Layout
KV blocks are allocated as uniform memory chunks. For a block size of $B_{\text{size}} = 16$ tokens:
$$\text{Stride}{\text{block}} = B{\text{size}} \times H_{KV} \times D \times \text{sizeof}(\text{dtype})$$
The metadata tracking engine manages these mappings using an explicit Block Table mapping logical sequence IDs to distributed physical pointers.
#include
#include
#include
#include
struct PhysicalBlock {
uint32_t device_id;
uintptr_t base_address;
uint32_t block_index;
uint32_t ref_count;
bool is_pinned_hbm;
};
class DistributedBlockTable {
public:
DistributedBlockTable(size_t block_size, size_t total_blocks)
: block_size_(block_size) {
free_blocks_.reserve(total_blocks);
for (size_t i = 0; i < total_blocks; ++i) {
free_blocks_.push_back(i);
}
}
bool AllocateSequenceBlocks(uint64_t seq_id, size_t num_tokens,
std::vector& allocated_blocks) {
size_t blocks_needed = (num_tokens + block_size_ - 1) / block_size_;
if (free_blocks_.size() < blocks_needed) {
return false; // Triggers Tier-2 Host DRAM eviction
}
for (size_t i = 0; i < blocks_needed; ++i) {
uint32_t block_idx = free_blocks_.back();
free_blocks_.pop_back();
PhysicalBlock block{
.device_id = 0,
.base_address = physical_pool_base_ + (block_idx * block_stride_),
.block_index = block_idx,
.ref_count = 1,
.is_pinned_hbm = true
};
allocated_blocks.push_back(block);
}
table_[seq_id] = allocated_blocks;
return true;
}
void FreeSequence(uint64_t seq_id) {
auto it = table_.find(seq_id);
if (it != table_.end()) {
for (const auto& block : it->second) {
free_blocks_.push_back(block.block_index);
}
table_.erase(it);
}
}
private:
size_t block_size_;
size_t block_stride_ = 16 * 8 * 128 * sizeof(uint16_t); // Example stride
uintptr_t physical_pool_base_ = 0x7f0000000000;
std::vector free_blocks_;
std::unordered_map> table_;
};
Chunked Prefill and Asynchronous Tensor Parallel Scheduling
Disaggregation removes decode interference from compute-dense nodes, but long-prompt prefill bursts still cause resource bubbles within the prefill cluster. If a sequence with $L=32768$ is processed monolithically, it blocks all other prefill executions on that TP rank for hundreds of milliseconds.
Chunked Prefill divides lengthy input contexts into discrete token chunks of size $C$ (e.g., $C = 512$ or $1024$). The engine co-schedules these prompt slices across execution cycles, interleaving prompt chunks from multiple requests or piggybacking small decode phases within residual compute budgets.
Un-chunked Execution (High Bubbles / Head-of-Line Blocking):
[ Request A: 8192 Tokens Prompt GEMM (Blocked for ~85ms) ]
[ Request B: 512 Tokens ]
Chunked Prefill Execution (Bounded Execution Budgets):
[ Req A Chunk 0: 1024 ] [ Req B Full: 512 ] [ Req A Chunk 1: 1024 ] [ Req A Chunk 2: 1024 ] ...
| <--- Slot 1 (12ms) -> | <-- Slot 2 (6ms) - | <--- Slot 3 (12ms) -> |
Mathematical Formulation of Chunked Attention
When processing chunk $k$ of a prompt spanning indices $[k \cdot C, (k+1) \cdot C - 1]$, the query tokens $Q_k$ must attend to:
- Keys and values within the current chunk: $K_{\text{local}}, V_{\text{local}}$ using causal masking.
- Keys and values computed during prior chunks: $K_{0:k-1}, V_{0:k-1}$ stored in the Paged KV-Cache without causal masking.
The attention computation updates the FlashAttention running log-sum-exp variables $(m, \ell)$ to merge partial softmax normalizations:
$$m_{\text{new}} = \max(m_{\text{prev}}, m_{\text{current}})$$
$$\ell_{\text{new}} = e^{m_{\text{prev}} - m_{\text{new}}} \ell_{\text{prev}} + e^{m_{\text{current}} - m_{\text{new}}} \ell_{\text{current}}$$
$$O_{\text{new}} = \text{diag}\left(e^{m_{\text{prev}} - m_{\text{new}}}\right) O_{\text{prev}} + \text{diag}\left(e^{m_{\text{current}} - m_{\text{new}}}\right) O_{\text{current}}$$
This online softmax update guarantees exact mathematical equivalence with monolithic attention while bounding single-kernel execution time to deterministic latency envelopes:
$$\tau_{\text{chunk}} \le \frac{2 \cdot C \cdot d_{\text{model}} \cdot \text{FLOPs}_{\text{layer}}}{\text{Peak Tensor Core Throughput}}$$
High-Performance Kernel Implementation: Triton Chunked Attention
The following Triton kernel demonstrates the compute path for a chunked prefill attention step. It loads historical KV cache blocks from a paged memory layout, processes the causal intra-chunk attention, and computes online softmax accumulation against preceding token contexts.
import torch
import triton
import triton.language as tl
@triton.jit
def _chunked_prefill_paged_kernel(
Q, K_Paged, V_Paged, Block_Tables, Output,
stride_qm, stride_qk,
stride_bk, stride_bn, stride_bd,
stride_om, stride_on,
sm_scale,
num_heads: tl.constexpr,
head_dim: tl.constexpr,
block_size: tl.constexpr,
chunk_size: tl.constexpr,
BLOCK_M: tl.constexpr,
BLOCK_N: tl.constexpr,
):
# Program identifiers
pid_m = tl.program_id(0)
pid_h = tl.program_id(1)
pid_b = tl.program_id(2)
# Offset calculations for queries
offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
offs_d = tl.arange(0, head_dim)
# Pointer arithmetic for Q
q_ptr = Q + pid_b * stride_qm + pid_h * head_dim + (offs_m[:, None] * stride_qk + offs_d[None, :])
q = tl.load(q_ptr, mask=offs_m[:, None] < chunk_size, other=0.0)
# Online Softmax accumulators
m_prev = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf")
l_prev = tl.zeros([BLOCK_M], dtype=tl.float32)
acc = tl.zeros([BLOCK_M, head_dim], dtype=tl.float32)
# Fetch physical block index from paged table
block_table_ptr = Block_Tables + pid_b * stride_bn
# Iterate across historical paged KV context
num_blocks = (chunk_size + block_size - 1) // block_size
for b_idx in range(0, num_blocks):
physical_block_id = tl.load(block_table_ptr + b_idx)
# Base pointers for physical paged K and V
k_page_ptr = K_Paged + physical_block_id * stride_bk + pid_h * head_dim
v_page_ptr = V_Paged + physical_block_id * stride_bk + pid_h * head_dim
offs_n = tl.arange(0, BLOCK_N)
k_ptr = k_page_ptr + (offs_n[None, :] * stride_bd + offs_d[:, None])
v_ptr = v_page_ptr + (offs_n[:, None] * stride_bd + offs_d[None, :])
# Load keys and values
k = tl.load(k_ptr, mask=offs_n[None, :] < block_size, other=0.0)
v = tl.load(v_ptr, mask=offs_n[:, None] < block_size, other=0.0)
# Compute QK^T
qk = tl.zeros([BLOCK_M, BLOCK_N], dtype=tl.float32)
qk += tl.dot(q, k) * sm_scale
# Causal mask application if operating on the diagonal boundary
col_indices = b_idx * block_size + offs_n
mask = offs_m[:, None] >= col_indices[None, :]
qk = tl.where(mask, qk, float("-inf"))
# Online softmax update cycle
m_curr = tl.maximum(m_prev, tl.max(qk, 1))
p = tl.exp(qk - m_curr[:, None])
l_curr = tl.exp(m_prev - m_curr) * l_prev + tl.sum(p, 1)
# Rescale accumulator
acc_scale = tl.exp(m_prev - m_curr)
acc = acc * acc_scale[:, None]
acc += tl.dot(p.to(v.dtype), v)
# Commit softmax tracking parameters
l_prev = l_curr
m_prev = m_curr
# Epilogue normalization
acc = acc / l_prev[:, None]
out_ptr = Output + pid_b * stride_om + pid_h * head_dim + (offs_m[:, None] * stride_on + offs_d[None, :])
tl.store(out_ptr, acc.to(Output.dtype.element_ty), mask=offs_m[:, None] < chunk_size)
Asynchronous Communication-Computation Overlap Pipeline
To prevent tensor migration over the RDMA network from stalling decode clusters, modern runtimes use an asynchronous multi-stream scheduling state machine.
Prefill GPU (P-Node) Decode GPU (D-Node)
Time --------------------- --------------------
| [ Chunk 0 Compute ]
| [ Chunk 1 Compute ] ----------> [ RDMA Ingress: Chunk 0 ]
| [ Chunk 2 Compute ] ----------> [ RDMA Ingress: Chunk 1 ]
| [ Epilogue / Done ] ----------> [ Finalize Block Table ]
v [ Step 0 Decode Compute ] (Immediate)
Double-Buffering Orchestration
By organizing transfer channels across dual-buffered CUDA streams, the memory movement of chunk $k$ occurs concurrently with the forward matrix multiply of chunk $k+1$.
import torch
class AsyncDisaggregatedPipeline:
def __init__(self, model_layers, num_streams=2):
self.compute_stream = torch.cuda.Stream()
self.comm_stream = torch.cuda.Stream()
self.model_layers = model_layers
# Dual-buffered staging allocations (Pinned Host / GPUDirect RDMA mapped)
self.transfer_buffers = [
torch.empty((1, 8, 512, 128), dtype=torch.float8_e4m3fn, device="cuda")
for _ in range(num_streams)
]
def forward_and_transfer_pipelined(self, prompt_chunks):
num_chunks = len(prompt_chunks)
for i in range(num_chunks):
# Compute current chunk on the compute stream
with torch.cuda.stream(self.compute_stream):
k_out, v_out = self.model_layers.forward_chunk(prompt_chunks[i])
buf_idx = i % 2
self.transfer_buffers[buf_idx].copy_(k_out)
# Synchronize transfer stream with compute completion of current chunk
self.comm_stream.wait_stream(self.compute_stream)
# Asynchronously dispatch buffer via RDMA over communication stream
with torch.cuda.stream(self.comm_stream):
self.issue_gpudirect_rdma_put(
source_tensor=self.transfer_buffers[buf_idx],
dest_node_rank=1,
remote_offset=i * 512
)
# Synchronize before handing off decode state
torch.cuda.synchronize()
def issue_gpudirect_rdma_put(self, source_tensor, dest_node_rank, remote_offset):
# Wraps low-level IBverbs / UCX rdma_write calls using external bindings
pass
Performance Engineering and End-to-End Latency Profiles
Evaluating the system dynamics of monolithic versus disaggregated architectures requires examining operational boundaries via the Roofline Model.
Log(Attainable Performance [TFLOP/s])
^
Peak |---------------------------------------+ (Compute Bound: Prefill Cluster)
Tensor Core /
FLOP/s /
| /
| /
| /
| / (Memory Bandwidth Bound: Decode Cluster)
| /
| / Slope = Peak HBM Bandwidth (TB/s)
| /
+---------------------+----------------------------------->
0.1 10 1000 Log(Operational Intensity)
Monolithic Bottleneck Analysis
In monolithic inference setups, when a prefill batch ($L=4096$) is interleaved with active decode streams, the compute engine toggles between the operational ceiling ($\approx 989 \text{ TFLOP/s}$ FP16 Tensor Core) and the sloped memory bandwidth bound ($3.35 \text{ TB/s}$).
Because the tensor cores and memory execution units share power budgets, instruction dispatch queues, and L2 cache crossbars, this switching incurs high overhead:
- L2 Cache Thrashing: Prefill GEMMs sweep large activation inputs through the $50\text{--}60 \text{ MB}$ L2 cache, completely evicting cached KV pointers required by concurrent decode queries.
- Frequency Throttling: Sustained high-power FP8/FP16 GEMMs cause dynamic core frequency scaling, throttling execution clock speeds for lighter memory-bound decode kernels.
Goodput Scaling Under Disaggregation
Disaggregation isolates these execution domains, optimizing system throughput within target latency limits (Goodput).
System Goodput (Requests/sec within Strict SLO Bounds)
^
| / Disaggregated Serving
| / (Stable TTFT & Flat ITL)
| /
| /
| /
| Monolithic Serving /
| (Tail Latency Collapse) /
| + /
| / \ /
| / \ /
| / \___________/
+----------+---------------------------------------------------->
0 100 500 1000 Concurrency
+------------------------------------+--------------------+--------------------+
| Metric | Monolithic Serving | Disaggregated VLLM |
+------------------------------------+--------------------+--------------------+
| Time to First Token (TTFT, P99) | 840 ms | 112 ms |
| Inter-Token Latency (ITL, P99) | 48.2 ms | 8.4 ms |
| ITL Normalized Jitter (Std. Dev.) | 21.6 ms | 0.9 ms |
| Maximum Attainable Batch Size | 128 seqs/node | 512 seqs/node |
| GPU SM Utilization Efficiency | 38.4% | 84.1% |
+------------------------------------+--------------------+--------------------+
Interconnect Bandwidth Sizing Guidelines
To architect a zero-stall disaggregated serving cluster, verify that network ingestion rate matches decode memory allocation capabilities:
$$\text{Required RDMA Bandwidth} \ge \frac{N_{\text{Prefill-Out}} \times \text{Size}(\text{KV}_{\text{Context}})}{\text{Target TTFT Budget} - \text{Prefill Execution Duration}}$$
Provisioning a 4:1 decode-to-prefill node ratio over dual-rail $400 \text{ Gbps}$ RoCEv2 switches avoids network queuing delays, keeping the disaggregated transport overhead within $5\text{--}8%$ of pure mathematical forward pass time.
Conclusion
Disaggregating LLM inference into decoupled, phase-specialized compute tiers addresses the structural limitations of monolithic continuous batching architectures. By assigning prompt prefill workloads to compute-dense clusters optimized for dense matrix mathematics, while routing sequential token generation to high-bandwidth, memory-dense decode nodes, disaggregated architectures resolve compute-memory resource interference.
Combined with Paged KV-Cache virtualization, chunked prompt scheduling, and pipelined RDMA memory transport, disaggregated inference transforms large-scale generative AI deployment. The resulting infrastructure delivers predictable tail latencies, sustained multi-node hardware saturation, and scalable throughput across enterprise production environments.
References
- PagedAttention & vLLM: Efficient Memory Management for Large Language Model Serving with PagedAttention
- DistServe Systems Research: DistServe: Disaggregating Prefill and Decoding for Goodput-optimized Large Language Model Serving
- FlashAttention Core Algorithms: FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning