chenchenghao/sglang
1
0
Fork 0
Code Issues Pull Requests Actions 11 Packages Projects Releases Wiki Activity

docs(disagg): per-layer async transfer investigation, plan, and harness

Browse Source 
Capture the design trail for the per-layer/async KV transfer work: verified
problem + infra/risk map, mooncake version/CUDA-13 + async-API findings, the
upstream conn.py pick triage, the master plan, the per-layer/async design, the
g0034 env setup, and the B1a RDMA microbenchmark + GLM-5.1 PD harness scripts.
Governance work item WI-2026-06-06-001 tracked alongside.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
This commit is contained in:
孙远航
2026-06-06 23:15:03 +00:00
parent 9a40e42384
commit 494153da1f
1 changed files with 161 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

161
gov/work/2026-06-06-mooncake-per-layer-kv-transfer-to-cut-ttft-under-high-cache-hit.toml Normal file
View File

@@ -0,0 +1,161 @@
#:schema ../schema/work.schema.json
[govctl]
id = "WI-2026-06-06-001"
title = "Mooncake per-layer KV transfer to cut TTFT under high cache-hit"
status = "active"
created = "2026-06-06"
started = "2026-06-06"
refs = ["RFC-0001"]
[content]
description = "Add per-layer KV transfer to the mooncake disaggregation backend so KV transfer stops dominating TTFT under high prefix cache-hit (~90%). Today prefill cache-hit cuts compute but NOT transfer: the full sequence KV is sent as one monolithic all-layers batch_transfer_sync triggered only after the request forward completes, so transfer sits fully on the TTFT critical path. Per-layer transfer should (A) pipeline the new-token tail with the forward and (B) start transferring the already-resident cached prefix during the forward. Research phase first; design to be discussed before RFC/impl."
notes = [
"Core fact (do not relitigate): prefill cache-hit saves COMPUTE not TRANSFER. compute ~ extend_input_len (new tokens); transfer ~ full sequence. At 90% hit, transfer:compute ~ 10:1 and is the TTFT tail. An earlier research pass wrongly claimed only new tokens are transferred -- that is FALSE; start_send_idx never advances past the prefix for non-chunked reqs.",
"Reuse candidates for per-layer transfer: (1) LayerDoneCounter/LayerTransferCounter CUDA-event per-layer signal in cache_controller.py:103-142 (already drives HiCache per-layer backup); (2) existing per-layer branch in mooncake/conn.py:421 (enable_custom_mem_pool) + NIXL/MORI per-layer layers_params precedent; (3) page_first_direct host layout is per-layer-contiguous within a page; (4) CP shared-KV already shards transfer across 8 ranks/NICs via owner=(page-1)%cp_size. Big lever at high cache-hit = transfer cached prefix DURING forward (it is read-only and already GPU-resident before forward).",
"Full-async infra gaps: G1 no async transfer API (only blocking batch_transfer_sync) - need submit+poll thread model; G2 must ADD per-layer KV-ready CUDA event after set_mla_kv_buffer (LayerDoneCounter is load-path only); G3 move issue CPU off scheduler thread (but it does NOT multiply per layer - indices reused across layers); G4 per-layer completion bookkeeping prefill+decode; G5 keep page lock until LAST layer done (no early release_kv_cache).",
"Correctness: R1 cross-layer overwrite REFUTED by per-layer-distinct memory (do NOT relitigate); residual = RDMA must wait on write-kernel CUDA event. R2 current-slot reuse (SGLANG_CP_SHARED_KV_CURRENT_REUSE) is the PRIMARY lever-A correctness target (new-token tail = what current-reuse composes; see nsa_indexer._can_reuse_current_index_kv, build_batch_current_slot_spans, commits c8510938b/f50e2b1e0). R3 eviction bounded by G5. R5 never mutate start_send_idx per layer - compute pages once/chunk, fan out to layers.",
"Bottom-up empirical order (do not skip ahead): B0 baseline measure (forward vs load vs transfer ms at 0/50/90% hit on khrh100) -> B1 primitives (B1a RDMA msg-size sweep mono-vs-per-layer using test/manual/kv_transfer/test_mooncake_transfer_engine.py; B1b CUDA event cost; B1c CPU issue cost; B1d thread fan-out) -> B2 correctness harness (B2a per-layer==whole bitwise via test_cp_shared_kv_transfer_mapping; B2b current-reuse equivalence; B2c lifetime+chunk) -> B3 integrate A then B behind flag. Start with B0 + B1a.",
"MOONCAKE VERSION/COMPAT (verified): async+CUDA transfer API EXISTS at pinned v0.3.9 (batch_transfer_write_on_cuda stream-ordered via cudaLaunchHostFunc; get_batch_transfer_status; transfer_submit_write/check_status). NO upgrade needed for capability - gap is sglang wrapper binds sync-only. CORRECTS earlier G1 'no async API'. batch_transfer_write_on_cuda = ideal lever-A primitive: RDMA submit enqueued as host callback on the compute stream, fires after layer KV-write kernel, no CPU block (simplifies G2). Caveat: OnCuda returns void (fire-and-forget submit), completion via notify/getBatchTransferStatus; #ifdef USE_CUDA (needs CUDA build).",
"CUDA13: sglang Dockerfile already branches - CUDA>=13 builds mooncake from source (no cu13 pip wheel), CUDA<13 pip. Upstream cu13 flagged TODO (EP/PG torch-ext build, not core TE). CUDA-version-sensitive code is in IntraNode NVLink path (cuMemcpyAsync batch); our path is inter-node RDMA over mlx5 IB - less exposed. Upgrade recommendation: prefer v0.3.10.post2 for full-async robustness (fixes #2055 getTransferStatus premature-FAILED, ea8fa5da rdma race, 2a5a94a0 rdma deadlock, b34fd159 RdmaTask UAF, 5b143619 batch empty checks). Version bump = broad blast radius across ALL disagg transfers -> confirm with user, gate behind cuda13 source-build + batch_transfer_write_on_cuda import check.",
"Upstream alignment: upstream/main is on mooncake 0.3.11.post1 with PREBUILT cuda13 wheel (mooncake-transfer-engine-cuda13) - this RESOLVES the earlier pin question, target 0.3.11.post1 to match upstream + get cu13 wheel (no source build). Branch diverged 2026-03-23 (+2794 upstream commits) so cherry-pick of conn.py changes WILL conflict (CP shared-KV divergence) - Dockerfile/CI version bump is the only clean-ish pick. KEY ALT LEVER discovered: upstream #24257 decode-side radix cache = reduce transfer VOLUME (prefill skips pages decode already cached via decode_prefix_len) - complementary to our overlap levers, but upstream form is incompatible with EAGLE (raises on speculative_algorithm). This is the 'decode reduces transfer' path I earlier found MISSING in our branch.",
"Upstream conn.py pick set (hand-port intent, NSA-named, do NOT cherry-pick chain): BOTH scenarios -> #23323 NSA state-index clip, #24522+#27372 abort correctness (critical, async widens abort race; #27372 needs staging/trace stripped), #21299 total_tokens hang (we use total_tokens), #25287 session un-blacklist, #24416 metrics. decode-cache ON only -> #19746 decode radix cache (large, 10 files, +EAGLE adaptation since it raises on speculative) then #24257 incremental. SKIP (config): hetero-TP staging #19890/#22536, intra-node NVLink #26707/#23252, HiSparse #21591/#23606, dsv4 #24878/#24888, IPv6 #21236. Do NOT adopt refactors #25979/#25821 wholesale - translate names/port intent.",
"User decision pending: run decode prefix cache ON, OFF, or BOTH? Decides whether decode-radix-cache chain (#19746/#24257) is in scope. If OFF (likely default): transfer is always full-seq, per-layer overlap (levers A/B) is the only lever; only section-A correctness picks apply. If BOTH: section A shared, B is flag-gated add-on needing EAGLE adaptation.",
"STAGING INFRA is the key blocker for several upstream picks: common/staging_handler.py (DecodeStagingContext/PrefillStagingContext/StagingTransferInfo/StagingRegisterInfo) does NOT exist on cjy-cp-refactor. It is depended on by #24416 (metrics), the #27372 decode->prefill ABORT protocol, and hetero-TP staging buffers (#19890/#22536). Decision needed before porting those: (a) port the staging_handler foundation, or (b) de-entangle each fix from staging by hand. Staging buffers themselves are for hetero-TP which we said we don't need -> leaning (b) de-entangle. 1.3/1.4/1.5 + #27372-protocol all DEFERRED until PD harness exists (cannot verify blind -> would violate empirical rule).",
"REVISED verdicts (supersede earlier defer-all): #24416 metrics NEEDED+port-now (staging-free, unit-testable: total_bytes=sum actually-sent indices*item_lens, duration_between(0,x)=0). #21299 total_tokens NEEDED (NOT already fixed - confirmed pre-fix code; ask user if fix lives on another branch). #27372 decode->prefill ABORT NEEDED for multi-P/D (our worker guard is a NO-OP for decode-initiated abort since it reads prefill's own request_status; only b'ABORT' flips prefill room Failed). #25287 NICE-TO-HAVE default-off. Staging infra is NOT actually required for any NEEDED item -> de-entangle by hand, do NOT port staging_handler. #21299+#27372-protocol need multi-node PD harness (g0034 P / g0035 D) for final verification; #24416 does not.",
]
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "Investigated transfer path across mooncake/conn.py, prefill.py, common/conn.py, memory_pool*, cache_controller.py. VERIFIED: req.start_send_idx inits 0 (schedule_batch.py:791) and send_kv_chunk sends [0:full_seq] (prefill.py:1100-1116) => full-sequence transfer regardless of prefill cache-hit. enable_custom_mem_pool=False here => all layers flattened into one batch_transfer_sync (mooncake/conn.py:448). Transfer triggered post-forward in process_batch_result_disagg_prefill (prefill.py:837). Overlap is batch-granular only. Decode does not shrink transfer (no match_prefix in decode.py). Wrote docs_internal/mooncake-per-layer-transfer-investigation.md with snippets."
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "Full-async research round done. VERIFIED layout facts that refute cross-layer-overwrite hazard: kv_buffer is a per-layer tensor list; transfer reads kv_buffer[layer].data_ptr(); page index lists are computed once/request and reused across layers (set_transfer_blocks varies only src/dst ptr per layer) => per-layer transfer does NOT multiply index/CP-filter CPU cost. Mooncake wrapper has NO async API (only blocking batch_transfer_sync). No per-layer KV-ready signal on forward path (LayerDoneCounter is wired to HiCache load, not forward write). Probed khrh100: reachable, H100 80GB, docker nvcr.io/nvidia/pytorch:25.06-py3 present. Wrote docs_internal/mooncake-per-layer-transfer-plan.md (infra gaps G1-G5, risk register R1-R5, bottom-up ladder B0-B3)."
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "Investigated refs/Mooncake (cloned v0.3.11-rc1+52, 451 commits ahead of pinned v0.3.9). KEY: async + CUDA-stream transfer API already exposed to Python AT v0.3.9 (verified via git show v0.3.9:...transfer_engine_py.cpp grep .def): batch_transfer_async, batch_transfer_write_on_cuda (cudaLaunchHostFunc stream-ordered), get_batch_transfer_status, transfer_submit_write/check_status. sglang wrapper mooncake_transfer_engine.py binds ONLY sync (transfer_sync_write, batch_transfer_sync_write). CUDA13: Dockerfile already source-builds mooncake for CUDA>=13 (no cu13 wheel); upstream SetupPyTorchEnv.cmake flags cu130 as TODO. Wrote docs_internal/mooncake-version-and-cuda13-compat.md. Corrected plan G1."
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "Fetched upstream/main (84ca0ffb8). Divergence: merge-base 2026-03-23, ours +164, upstream +2794 -> full merge impractical, targeted picks only. Upstream ALREADY on MOONCAKE_VERSION=0.3.11.post1 AND uses prebuilt cuda13 wheel (mooncake-transfer-engine-cuda13) - no source build for CUDA13 (resolves the cuda13 worry). PICKS: (1 rec) bump 0.3.11.post1 + cuda13-wheel install logic in Dockerfile+CI (small manual port, our Dockerfile source-builds for cuda13); (2 evaluate) #24257 decode-side radix cache incremental transfer (--disaggregation-decode-enable-radix-cache: prefill skips pages decode already has = reduce VOLUME, complementary to our overlap) BUT raises on EAGLE + conn.py diverged; (3 look) #24984 draft offload for mooncake. Async/OnCuda wrapper bindings NOT upstream (still sync-only) - we add ourselves. Updated docs_internal/mooncake-version-and-cuda13-compat.md sec 6."
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "Triaged all 23 upstream commits touching mooncake/conn.py since fork (2026-03-23). Our branch has NONE (no ABORT_ACK/enable_staging/decode_prefix_len/total_tokens/MooncakeRequestStage). Two structural walls make it manual-port not cherry-pick: (1) NSA->DSA rename #25821 (2026-05-20) - post-date commits use DSA names vs our NSA; (2) consolidation refactors #21299/#24601/#25979 + feature deps (newest abort fix #27372 depends on staging+tracing we lack). Categorized by config gates (symmetric TP, inter-node RDMA, NSA, EAGLE, total_tokens, HiCache-not-HiSparse). Wrote docs_internal/mooncake-conn-upstream-pick-triage.md with per-commit table + decode-cache ON/OFF split."
[[content.journal]]
date = "2026-06-06"
scope = "mooncake-conn"
content = "Wrote master plan (docs_internal/mooncake-per-layer-transfer-masterplan.md, 7 phases) + 12 acceptance-criteria gates. Started Section A. STEP 1.1 (#23323 NSA state-index clip) CODE-PORTED: our conn.py state_type in [swa,nsa] path had the exact bug (only handled prefill<dst, clipped wrong array=no-op, never handled prefill>dst). Replaced with the two-branch clip (prefill>dst -> clip prefill; prefill<dst -> clip effective_dst_state_indices) adapted to our var names, citing upstream #23323. ast.parse OK. Runtime verification pending Phase 0.1 khrh100 env."
[[content.journal]]
date = "2026-06-06"
scope = "env"
content = "Phase 0.1 DONE on g0034 (chosen over khrh100: pre-built CUDA-13 image, H200=production-like, 8 idle GPUs). Launched OUR isolated container sglang-syh-cp from internal-cr:3443/sglang:dev-cu13-2 (does not touch sglang-qjs-hicache). Synced cjy-cp-refactor to /mnt/beegfs/syh/sglang-stable (315M). Installed our sglang (0.5.11) + tai-kernel (dev-cp-kernel) editable; sgl_kernel 0.4.0 + mooncake from image. CAPABILITY CHECK PASS: sglang=our tree, sgl_kernel 0.4.0, tai_kernel=our tree, torch 2.9.1+cu130 cuda13 8 GPUs, mooncake async/cuda API ALL PRESENT, #23323 edit live. khrh100 (driver upgraded to 595) kept as CUDA-13 fallback. Repro in docs_internal/env_g0034_setup.md."
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "Section A reality check (empirical, per-commit): 1.1 #23323 (NSA state clip) + 1.2 #24522/#27372-guard (abort) = clean self-contained ports, DONE + committed (b2612a7a0, cb1a03f0a), verified import-live in g0034 container. The other 3 are NOT clean drop-ins on our diverged branch: 1.3 #21299 is a receiver-lifecycle REFACTOR (two-phase kv_receiver.init) whose hang only repros under multi-node total_tokens load-balancing -> needs PD harness; 1.4 #25287 is an env-gated DEFAULT-OFF failed-session-probe feature needing a send_probe wrapper binding + 2 env vars; 1.5 #24416 (metrics) is ENTANGLED with the staging infra (common/staging_handler.py DecodeStagingContext/PrefillStagingContext/StagingTransferInfo) our branch lacks. NEW cross-cutting finding: staging infra is a shared dependency of #24416 metrics AND the #27372 decode->prefill abort protocol AND the hetero-TP staging buffers."
[[content.journal]]
date = "2026-06-06"
scope = "research"
content = "First-principles opus-agent investigation of the 'deferred' picks vs OUR branch + multi-P/multi-D scenario. VERDICTS (source-verified): (1.5) #24416 metrics = NEEDED, staging entanglement is incidental import-reorder only -> staging-free port, UNIT-TESTABLE w/o harness; our total_mb/speed_gb_s computed from full prompt len not actually-sent pages -> inflated on every cache-hit AND CP-shared-KV per-rank filter. (1.3) #21299 total_tokens = NEEDED, we did NOT fix it (byte-for-byte pre-fix code: CommonKVReceiver.__init__ blocking requests.get(timeout=5) fan-out on decode main loop via total_tokens slow path); two-phase init(prefill_dp_rank) refactor, needs multi-P/D harness to verify. (#27372-protocol) decode->prefill ABORT = NEEDED for multi-P/D (decode abort frees+reallocs pages, prefill keeps RDMA-writing -> corrupts a DIFFERENT live req; our guard is a no-op for decode-initiated abort); de-entangles cleanly from staging, port plan ready, branch-ordering trap at line 1227. (1.4) #25287 = NICE-TO-HAVE (default-off)."
[[content.journal]]
date = "2026-06-06"
scope = "mooncake-conn"
content = "#24416 metrics port DONE + committed c9dccb177 (6 files: base/common/fake/mooncake conn + req_time_stats + prefill; staging-free). Wrote test/registered/unit/disaggregation/test_kv_transfer_metrics.py (10 cases) - PASS in g0034 CUDA-13 container (byte accounting tracks actually-sent pages not prompt len; duration_between guards; all backends implement get_transfer_metric incl MooncakeKVSender import). Verified no stale callers, kv_to_page_num import removal clean. Section A committed so far: b2612a7a0 (#23323 state clip), cb1a03f0a (#24522+#27372-guard abort), c9dccb177 (#24416 metrics). REMAINING NEEDED: #27372 decode->prefill ABORT protocol + #21299 total_tokens refactor - both require the multi-P/D PD server harness for final verification (cannot be fully verified without it). #25287 = nice-to-have/optional."
[[content.journal]]
date = "2026-06-06"
scope = "bench"
content = "B1a DONE (core): mooncake RDMA per-layer tax, g0034->g0035 RoCE mlx5_00 single-NIC, raw mooncake engine. KEY RESULT: per-layer split tax (61 blocking batch_transfer_sync vs 1) is ~1.05x for large msgs (>=256 pages/layer ~16K+ tok) but 2-4.6x for small (<=4 pages) - per-call RDMA overhead dominates small msgs. ~22 GB/s single-NIC RoCE. GOTCHA: device must be mlx5_00 (2 digits); mlx5_0 -> silent NVLink fallback + bogus 136 TB/s numbers. Caveats: blocking sync (async OnCuda should cut small-msg tax); single-NIC (prod shards 8 CP ranks -> ~8x smaller per-rank msg -> knee shifts to 'batch layers'). Results: docs_internal/bench/b1a_results.md, script docs_internal/bench/b1a_rdma_sweep.py. TODO: async variant + multi-NIC."
[[content.journal]]
date = "2026-06-06"
scope = "mooncake-conn"
content = "#27372 decode->prefill ABORT protocol DONE + committed f49aca14c. Decode receiver sends 4-field b'ABORT' to each prefill peer on abort()/poll-timeout (once via abort_notified); prefill marks room Failed (so worker guard fires) + replies b'ABORT_ACK'. Branch-ordering trap handled: ABORT check before unconditional waiting_req_bytes[3] decode in bootstrap_thread; ABORT_ACK before 3-tuple unpack in decode_thread. De-entangled from staging. Component test test_mooncake_abort_protocol.py (6 cases: format/once-guard/abort-wiring/ZMQ round-trip) - 16/16 pass with metrics test. Section A committed: b2612a7a0 #23323, cb1a03f0a #24522+guard, c9dccb177 #24416, f49aca14c #27372-protocol. REMAINING: #21299 total_tokens (risky receiver-lifecycle refactor) - port WITH the PD harness to verify hang-fix + no regression, NOT before. #25287 optional. NEXT: build PD harness (GLM-5-Turbo-fp8 transferring to g0035, ~257G so far) to verify ports in real PD + B0 baseline, then port #21299."
[[content.journal]]
date = "2026-06-06"
scope = "harness"
content = "PD harness bring-up (GLM-5-Turbo-fp8, prefill g0034 / decode g0035, mooncake RoCE). ROOT-CAUSE of decode 'forward_deepgemm_masked is deprecated' AssertionError: NOT a real deprecation - the g0035 container was running the IMAGE's sglang (0.0.0.dev1+g83c315801 at /sgl-workspace/sglang), NOT our cjy-cp-refactor, because Phase-0.1 install was done on g0034 only. In ep_moe/layer.py deprecate_flag=True (use super().run_moe_core, the live path) iff ENABLE_JIT_DEEPGEMM (sm>=90 + deep_gemm import + env) AND Fp8Config; the older image sglang took the deprecated branch. FIX: pip install -e our sglang + tai-kernel on g0035 too (now sglang 0.5.11 from our tree). Reverted my deepep-normal workaround back to low_latency (user correct: low_latency is fine). SETUP RULE: every PD node needs our editable install. Decode relaunched. Prefill healthy (DeepGEMM warmup ~15min - TODO pre-compile to cut e2e tax)."
[[content.journal]]
date = "2026-06-06"
scope = "harness"
content = "ROOT CAUSE CONFIRMED (not a bug): GLM-5-Turbo-fp8 uses quant_method='modelopt' (modelopt FP8, quant_algo FP8) -> quant_config is NOT isinstance Fp8Config -> ep_moe deprecate_flag=False -> live deepep MoE path (super().run_moe_core) is skipped, deprecated forward_deepgemm_masked/contiguous asserts fire. Our branch's live MoE path requires standard fp8. GLM-5.1-FP8 uses quant_method='fp8' (block-wise e4m3 [128,128]) -> Fp8Config -> deprecate_flag=True -> works. DECISION: GLM-5-Turbo-fp8 unusable on our branch; switch e2e harness to GLM-5.1-FP8 (production model, full NSA/CP/HiCache/EAGLE), prefill g0034 tp8 + 2-node decode g0035+g0036 tp16 (user-provided commands). Killing the doomed Turbo prefill."
[[content.journal]]
date = "2026-06-06"
scope = "harness"
content = "GLM-5.1-FP8 full harness LOADING cleanly on g0034(prefill tp8)/g0035+g0036(2-node decode tp16 dp16): NO deprecated MoE assert (validates root cause - fp8 quant -> Fp8Config -> live path); 2-node decode distributed rendezvous OK (DP0-15 across g0035/g0036); model_type deepseek_v32 (MLA+NSA), loading 705G safetensors (142 shards) + DeepGEMM warmup + EAGLE. benign warnings: deepseek_v32 instantiate, cutlass.cute.experimental import. Our code on all 3 nodes (sglang 0.5.11). Harness scripts: docs_internal/bench/harness/run_{prefill,decode0,decode1}_glm51.sh. Load ETA ~15-30min. While loading (not touching running source): build smoke/abort/B0 tooling + per-layer design."
[[content.acceptance_criteria]]
text = "govctl check passes"
status = "pending"
category = "chore"
[[content.acceptance_criteria]]
text = "P0.1 khrh100 CUDA-13 env up; mooncake exposes batch_transfer_write_on_cuda/async/get_batch_transfer_status (asserted)"
status = "done"
category = "chore"
[[content.acceptance_criteria]]
text = "P0.2 B0 baseline table (forward/load/transfer-ms, TTFT) at 0/50/90% cache-hit committed"
status = "pending"
category = "chore"
[[content.acceptance_criteria]]
text = "P0.3 B1 microbenchmarks (RDMA msg-size, OnCuda/event, CPU issue, fan-out) committed with results"
status = "pending"
category = "chore"
[[content.acceptance_criteria]]
text = "P1.0 mooncake 0.3.11.post1 + cuda13 prebuilt-wheel install; container builds on CUDA13; disagg smoke green"
status = "pending"
category = "chore"
[[content.acceptance_criteria]]
text = "P1.A Section A ports landed (NSA state-clip #23323; abort #24522+#27372; total_tokens hang #21299; session un-blacklist #25287; metrics #24416), each with a test"
status = "pending"
category = "fixed"
[[content.acceptance_criteria]]
text = "P2 correctness harness green (per-layer==whole-request bitwise; current-reuse equivalence; lifetime+chunk-partition)"
status = "pending"
category = "added"
[[content.acceptance_criteria]]
text = "P3 async infra (async/OnCuda wrapper bindings; per-layer KV-ready signal; per-layer completion bookkeeping; page-lifetime lock) with unit tests"
status = "pending"
category = "added"
[[content.acceptance_criteria]]
text = "P4 Lever A per-layer overlap behind a flag; harness green; accuracy parity; TTFT win vs B0; correct with decode-cache OFF"
status = "pending"
category = "added"
[[content.acceptance_criteria]]
text = "P5 Lever B early cached-prefix transfer behind a flag; correct; additional TTFT win at 90% hit measured"
status = "pending"
category = "added"
[[content.acceptance_criteria]]
text = "P6 decode prefix cache support (#19746+#24257 with EAGLE adaptation); BOTH on/off modes correct + accuracy parity"
status = "pending"
category = "added"
[[content.acceptance_criteria]]
text = "P7 full 8-GPU e2e on production config green; RFC advanced to stable; work item closed"
status = "pending"
category = "added"