perf(maru_vllm): eliminate direct connector bottlenecks

[youngrok-XCENA]

🤔 Background & Motivation (Why)

MaruKVConnector의 control/data path에는 long-context에서 증폭되는 구조적 병목이 있었습니다.

• scheduler와 worker가 (layer × chunk) 단위로 단건 RPC와 작은 GPU 전송을 반복했습니다. 64k prompt에서는 요청당 key가 약 8,000개까지 늘어났습니다.
• chunk key를 prefix마다 다시 해싱해 scheduler lookup이 O(n²)으로 증가했습니다.
• chunked-prefill 경계에서 누적 block mapping을 잃어 일부 chunk가 누락되거나 잘못된 GPU slot을 저장할 수 있었습니다.
• worker의 대용량 CXL mmap/prefault/CUDA registration이 첫 populate 요청까지 지연되어 첫 요청과 첫 inflight wave를 수초 동안 멈췄습니다.

이 PR은 해당 병목과 정합성 문제를 제거하고, single-instance 및 P2P direct connector 검증 예제를 함께 제공합니다.

위 재설계로 캐시 히트 지표는 LMCache 경유(non-MP)와 동등해졌지만, 로드·저장이 여전히 엔진 안의 동기 호출이라 한 요청의 KV 전송이 같은 배치의 다른 요청 decode를 세우는 병목이 남아 있었습니다. 같은 조건에서 로드·저장을 별도 프로세스로 빼는 LMCache MP 모드는 동시 8요청 캐시 히트 TTFT가 29% 낮고 처리량이 19% 높았습니다. 후속 커밋 4개는 이 갭을 in-process에서 해소합니다 — 벌크 전송을 GPU copy engine으로 빼고, 전송 시점을 첫 토큰의 임계 경로 밖으로 옮깁니다.

🏗️ Design Changes

flowchart LR
    subgraph before["Before"]
        S1[Scheduler] -->|key별 exists RPC| M1[MaruServer]
        W1[Worker] -->|layer × chunk별 retrieve/store| M1
        W1 -->|첫 요청에서 CXL mapping| C1[CXL]
    end
    subgraph after["After"]
        S2[Scheduler] -->|batch exists and linear key hashing| M2[MaruServer]
        W2[Worker] -->|chunk-packed batch retrieve/store| M2
        W2 -->|multi-layer coalesced transfer| C2[CXL]
        I2[Engine startup] -->|eager worker connect| W2
    end
    before ~~~ after

Loading

• metadata lookup/store를 batch RPC로 묶습니다.
• 한 chunk의 모든 layer를 하나의 packed CXL slab/object로 저장합니다.
• Flash KV layout에서는 pinned CXL slab과 paged GPU cache 사이를 multi-layer kernel로 직접 전송하고, 지원되지 않는 환경에서는 기존 per-layer 경로로 fallback합니다.
• worker CXL mapping은 KV geometry가 확정되는 register_kv_caches()에서 수행해 요청 critical path에서 제거합니다.
• scheduler/worker role과 key scheme은 layerwise compatibility mode에서도 일치하도록 유지합니다.

후속 커밋(비동기화)의 구조 변화:

flowchart LR
    subgraph sync_path["Before — 엔진 안 동기 로드/저장"]
        direction TB
        A1["cache-hit/populate 전송이<br/>엔진 스레드에서 동기 실행 (SM 점유)"] --> A2["같은 배치의 다른 요청<br/>decode 정지"]
    end
    subgraph async_path["After — 비동기화 + 요청 내 겹침 (opt-in)"]
        direction TB
        B1["로드: 요청 파킹 후 로더 스레드가<br/>copy engine DMA로 수행"] --> B2["저장: prompt 완료 뒤<br/>write-behind로 발행"] --> B3["단독 hit 요청:<br/>layer k 연산 ∥ layer k+1 로드"]
    end
    sync_path ~~~ async_path

Loading

• 로드 비동기화 — 캐시 적중 요청을 vLLM의 비동기 로드 계약으로 잠시 세워두고, 백그라운드 로더가 조회와 CXL→GPU 전송을 수행합니다. 벌크 바이트는 copy engine이 옮기고 GPU 연산 유닛은 짧은 산포 단계만 씁니다. 스레드 분리만으로는 효과가 없었고(전송이 연산 유닛을 점유), copy engine으로 옮긴 뒤에야 decode 보호가 발동했습니다.
• 저장 write-behind — chunked prefill 중간 저장을 없애고 prompt가 완성된 뒤 full prefix를 한 번에 발행합니다. GPU staging을 거쳐 copy engine이 CXL로 내리고, 등록은 백그라운드에서 완료됩니다.
• layer 단위 겹침 — 진행 중인 적중 요청이 1건일 때는 요청 내부에서 layer k 연산과 layer k+1 로드를 겹칩니다. 프로세스 분리 구조(MP)에는 없는 축이라, 저동시성 TTFT에서 MP를 앞서는 원천입니다.
• 세 기능 모두 opt-in이며 기본값 off — 켜지 않으면 기존 동작과 동일합니다.

성능 개선 결과 (Llama-3.1-8B, lveval 16k, 문서 20개, CXL-DRAM)

지표	개선 전	개선 후	LMCache non-MP	LMCache MP
캐시 히트 TTFT (동시 1)	418 ms	124 ms (−70%)	136 ms	147 ms
캐시 히트 TTFT (동시 4)	—	249 ms	296 ms	262 ms
캐시 히트 TTFT (동시 8)	2,288 ms	462 ms (−80%)	648 ms	443 ms
캐시 히트 TPOT (동시 8)	52.9 ms	21.7 ms (−59%)	24.7 ms	22.0 ms
TPS (동시 8)	76.4	228.7 (3.0×)	188.7	229.7
populate TTFT (동시 8)	6,906 ms	2,549 ms (−63%)	3,737 ms	2,704 ms

• 개선 후는 이 PR 전체 적용(opt-in 플래그 3종 활성) 기준입니다. LMCache non-MP를 전 지표에서 넘고, LMCache MP와는 처리량/TPOT 동등·동시 1/4 TTFT와 populate 우위이며, 동시 8 TTFT +4%(462 vs 443 ms)만 남습니다.
• 개선 전은 2026-07-14 측정(동시 4 셀 없음), 개선 후·참조 열은 2026-07-22 같은 grid의 인접 run — run 간 TTFT 수% 드리프트를 감안해 주세요. 캐시 히트는 2 round 평균입니다.

📝 Implementation Details

• batch_retrieve/batch_store와 payload-bounded retrieve batching
• KV geometry 기반 CXL page 크기 자동 산정
• chunk-packed storage와 completion-marker 제거
• async layer load, optional deferred/fused paths, packed load/store stream 관리
• incremental blake2b chunk hashing 및 요청 단위 key memoization
• chunked-prefill의 누적 block IDs와 절대 chunk offset 보존
• batch-store 실패 시 CXL handle 정리 및 중복 prefix store 방지
• eager worker connect 실패 시 첫 실제 요청의 즉시 retry 동작 보존
• examples/vllm/single/ 자동 cold→warm 검증과 examples/vllm/p2p_sharing/ 예제 정리

lmcache.c_ops가 설치된 경우 multi-layer CUDA kernel을 선택적으로 재사용하지만 LMCache connector/service를 실행하지는 않습니다. 해당 extension이 없으면 호환 가능한 per-layer fallback을 사용합니다.

후속 커밋(비동기화):

• deferred load (maru_enable_deferred_loading) — scheduler가 hit 요청을 WAITING_FOR_REMOTE_KVS로 파킹, 단일 로더 스레드가 batch_retrieve 후 registered CXL mmap → GPU staging cudaMemcpyAsync(copy engine DMA) + multi_layer_kv_transfer 산포, 전용 stream의 CUDA event를 get_finished()가 폴링. 로드 실패 블록은 재계산으로 보고
• write-behind (maru_enable_write_behind) — save hook은 마지막 layer에서 metadata만 기록하고 prompt-완료 step의 get_finished()에서 발행. pending key와 shared-prefix waiter를 lock 아래 추적해 중복 저장 방지, 완료 전 끝난 요청은 connector가 GPU block ownership을 보류 후 finished_sending으로 해제, preemption은 store stream drain 후 진행
• layerwise overlap (maru_enable_layerwise_overlap, deferred 전제) — live deferred-hit 요청이 정확히 1건일 때만 packed layout을 유지한 pitched cudaMemcpy2DAsync + layer별 CUDA event를 발행하고 attention은 자기 layer event만 대기. 2건 이상이면 whole-request DMA 경로 유지. step-local 판정은 full grid에서 고동시성 TPOT 회귀(+16~18%)를 일으켜 live-request 추적으로 교체
• CPU/non-Flash fallback은 동기 store를 유지하고, CUDA 경로 실패 시 silent fallback 없이 해당 요청 블록을 invalid로 보고해 재계산

✅ Tests

• Unit tests — 801 passed, 1 skipped (후속 커밋 반영; connector 단독 99 passed. suite 순서에서만 나타나는 기존 CUDA host-memory 재등록 이슈 1건은 단독 실행 시 pass)
• Integration tests — Naru runner 16k-token, 20 documents, inflight 1/4/8 grid에서 기능별 A/B 검증; backend당 180/180 성공, external hit 66%(만점), 신규 store/load/CUDA 실패 0
• Manual tests — direct cache-hit output matched the reference path 20/20 byte-for-byte (동기 경로 기준; 후속 기능은 기본 off)
• Static checks — Ruff, format check, connector MyPy(기존 대비 신규 오류 0), example shell bash -n

장비 검증 결과 warm cache-hit TTFT/TPOT/TPS의 기존 병목이 제거됐고, populate TTFT도 if1 916 ms, if8 3,035 ms로 안정화됐습니다.

후속 커밋(모두 opt-in) 활성 시, 같은 grid에서 warm cache-hit TTFT는 동시 8에서 580→462 ms(−20%), 처리량 192→229 tok/s(+19%), populate TTFT는 동시 8에서 3,049→2,549 ms(−16.4%)입니다. 같은 조건의 LMCache MP 모드와 처리량/TPOT가 동등하고 동시 1/4 TTFT는 앞섭니다.

🔗 Related Issues (optional)

🌿 Related PRs (optional)

• Builds on fix: store all chunked-prefill steps in MaruKVConnector, not just the first #57 (chunked-prefill continuation store)

🌿 Related Branches (optional)

• None

📦 Release Note (for auto-generation / write in English)

NEW

CHANGED

FIXED

IMPORTANT NOTES

[github-actions]

Coverage Report

File	Stmts	Miss	Cover	Missing
__init__.py	6	0	100%
__main__.py	3	3	0%	5, 7–8
allocation_manager.py	129	11	91%	30–31, 44–45, 52, 237–238, 268–271
client.py	187	8	96%	88, 130, 147–148, 311–313, 319
config.py	46	1	98%	72
constants.py	9	0	100%
device_scanner.py	115	5	96%	13–15, 63, 96
handler.py	557	86	85%	153, 168, 179–186, 194–196, 211, 217–219, 227, 238–243, 248, 273, 300–301, 305, 345, 349–350, 355, 377, 384–385, 389–395, 398, 402–404, 420, 422–423, 432, 486, 507, 517–518, 624–626, 633, 755–758, 761, 772–775, 781, 867, 971, 1134–1138, 1144, 1155–1159, 1165, 1244, 1249, 1291
ipc.py	275	2	99%	365, 441
kv_manager.py	108	0	100%
logging_setup.py	19	0	100%
plugin.py	59	3	95%	65–67
protocol.py	233	0	100%
resource_manager_installer.py	103	13	87%	80–86, 167, 169–172, 187
rpc_async_client.py	190	0	100%
rpc_async_server.py	111	0	100%
rpc_client.py	71	0	100%
rpc_client_base.py	106	10	91%	185, 221–222, 233–234, 306–307, 311–312, 375
rpc_handler_mixin.py	106	19	82%	154–156, 159–161, 219–222, 227, 231–235, 256–257, 263
rpc_server.py	64	0	100%
serializer.py	81	0	100%
server.py	170	13	92%	64–65, 73, 168, 172, 243, 247, 265–266, 357–359, 444
stats_manager.py	95	0	100%
types.py	60	1	98%	145
uds_helpers.py	13	0	100%
memory
__init__.py	5	0	100%
allocator.py	55	0	100%
mapper.py	164	17	90%	65–74, 77–80, 293, 342, 376
owned_region_manager.py	101	1	99%	212
types.py	63	0	100%
TOTAL	3325	193	94%

Tests	Skipped	Failures	Errors	Time
724	4 💤	0 ❌	0 🔥	7.601s ⏱️