LMCache KV Transfer Pipeline — Internal Architecture

1. Root Cause: Why LMCache Is Slower

Section §2 established that LMCache adds +2.73 s per job relative to FCFS under the benchmark conditions (H200, JPS=1, Exp C). This section traces that overhead to a single, quantifiable mechanism: excessive CUDA stream synchronization inside wait_for_save, which blocks the vLLM scheduler main loop on every prefill step.

1.1 The Cost: wait_for_save Blocks the Scheduler

vLLM's KV Connector API requires the scheduler to call wait_for_save before dispatching each step. NVTX profiling of 51 completed jobs (H200, JPS=1, Exp C) captured 1,769 invocations. The distribution is strongly bimodal: a small minority of prefill-step calls account for almost all wallclock time.

1.2 Inside wait_for_save: cudaStreamSynchronize

cudaStreamSynchronize(stream) is a CUDA Runtime API call that blocks the calling CPU thread until every operation previously queued on the given stream has completed on the GPU. Internally the call performs a user-to-kernel transition (via ioctl to the NVIDIA kernel driver), checks stream completion status, sleeps if the GPU is still working (using futex or driver poll), and resumes on GPU interrupt. Even when the GPU finishes its DMA in 0.67 ms (measured via NVTX from_gpu range, see §1.3), the round-trip through the kernel — transition + interrupt latency + wakeup — adds ~3–4 ms of fixed overhead per call on PACE nodes.

The Syncs/call and Memcpy/call columns above link directly to the CUDA Runtime totals in the table below (3,744 = 159 × ~24; 2,996 = 159 × ~19).

The derivation table below shows that the measured 115.71 ms per prefill step is almost entirely explained by 24 sync calls × 4.74 ms each, leaving less than 2% unexplained residual.

Root Cause (corrected)

CUDA Stream Definitions

Overhead Decomposition (115.71 ms per prefill step)

1.3 The Hardware Floor: DMA Is Not the Bottleneck

The GPU-side DMA transfer (MEMCPY D2H) was measured separately via Nsight Systems GPU-activity records across 4 GPU families and 3 JPS values (12 measurements, Exp C). These measurements establish the minimum achievable time for the actual data movement, independent of CPU-side synchronization overhead.

On the H200, GPU-side DMA completes in 0.674 ms. The per-call cudaStreamSynchronize cost is 4.74 ms. Therefore the DMA itself accounts for only 0.674 ÷ 4.74 = 14.2% of each sync call; the remaining 85.8% is fixed kernel-mode overhead. The table below itemises this split.

1.4 Optimization Path: Why Sync Patches Failed

Two patches targeting synchronization patterns (Exp D and Exp F) were tested to validate the original (incorrect) root cause hypothesis. Both produced negative results — not because the patches were poorly implemented, but because the ~15 ms real offload work is actual GPU computation, not a sync overhead artifact that can be removed by rearranging sync calls.

Exp D — Coalesce store_stream Sync (Negative Result)

Removed per-chunk store_stream.synchronize() in gpu_connector.py:from_gpu(), retaining only the single sync at the end of batched_from_gpu(). Tested on H200, H100, A100 (same workload as Exp A, JPS=1 and JPS=6, control vs patched sequentially).

Exp F — stream.query() Guard (Negative Result)

Added a store_stream.query() guard before each store_stream.synchronize() in gpu_connector.py, so the sync is skipped when the stream is already idle. Results were also negative: JPS=1 latency increased on both GPUs tested.

2. Ablation Experimental Design

2.1 Instrumentation

2.2 Sync Ablation (4 Modes)

The initial hypothesis (§2) was that torch.cuda.synchronize() in LMCache's KV-onload path blocks the CPU scheduler thread and starves the GPU. To falsify or confirm this, we need to replace the synchronization primitive itself with progressively looser variants and see whether wall-clock latency improves. If it does, the sync primitive was the bottleneck; if it doesn't, the bottleneck is somewhere else. The tool is our_experiments/scripts/patch_sync.py, which runtime-patches gpu_connector.py inside the LMCache installation in the conda env (so vLLM picks up the change on server restart).

The target line in original LMCache (gpu_connector.py) is at the end of batched_to_gpu, after firing the async CPU→GPU KV load on a dedicated load_stream:

# NOTE: wait until both compute and load streams are done
torch.cuda.synchronize()   # ← the line under test

Each mode replaces exactly this call with a different primitive (the rest of LMCache code is untouched). All four modes are tested head-to-head against the same FCFS/Continuum baselines under identical JPS and workload. Here is what each mode does and what it is supposed to test:

self.load_stream.synchronize()
self.compute_stream.synchronize()

self._sync_event = torch.cuda.Event()
self._sync_event.record()
self._sync_event.synchronize()

pass  # sync removed for ceiling measurement

2.3 Transfer Ablation

2.4 Parameters

2.5 GPUs Tested

3. Ablation Results

3.1 Wall-Clock Performance (H200)

Log-scale comparison of all 7 policies. Continuum (green) is the only variant whose curve diverges downward at JPS≥6. The five LMCache variants cluster together above FCFS, confirming that no sync strategy rescues LMCache — see Finding 1 in §5.

JPS=10 Mean JCT Comparison

3.2 Scheduler Step Analysis (H200)

3.3 L40S Results (48GB VRAM) — where the mechanism differs

L40S shows a much more severe LMCache penalty than H200 — average batch size drops from 26.7 (FCFS) to 8.0 (LMCache default), a 3.3× reduction. On H200 the same comparison is 16.6 vs 17.6, basically unchanged. The obvious-sounding explanation — "L40S has less VRAM so LMCache's buffers eat into KV budget" — is wrong. The actual mechanism is a subtler interaction between vLLM's auto-tuned max_num_batched_tokens and LMCache's wait_for_save stall.

4. Deep Dive: KV Transfer Pipeline

4.1 vLLM Step Lifecycle

A vLLM step = one schedule() + one model.forward() + sampling. Each step produces one output token per request in the batch.

4.2 Per-JPS Breakdown: Why High Concurrency Hurts More

4.3 Full Call Chain

4.4 Why the blocking wait exists — upstream rationale

That sentence is the entire upstream design discussion. There is no GitHub issue debating the trade-off, no PR proposing a non-blocking alternative, no blog post elaborating on the constraint, and no design doc. The wait-for-save barrier is treated as a self-evident safety requirement rather than an explicit trade-off.

The paged block reuse hazard

The blocking wait_for_save forbids this by freezing the engine loop until every in-flight save has finished reading every dirty block. It is a correctness barrier for the block allocator, not a performance choice — there is no way to skip it without changing the semantics of the free pool.

Potential mitigations (none implemented upstream as of 2026-04)

5. Key Findings

6. Future Directions

7. Code & Data Files

Code Changes

Experiment Scripts

Results Directory