docs(kvc): freeze reseed slow-path audit + three reviewer challenges

Standalone reference document capturing the v2 reseed slow-path forensic
audit before opening the feat/d-to-p-sync branch. Designed to be quoted
directly by future paper drafts and to prevent the team from re-relitigating
the same questions verbally.

Contents:

§1. The three team-member challenges that disproved "capacity-backup will
    save the slow path" (each with code citation and verdict):
    1) P pool can't fit all backups -- replay.py:1618-1620 caps backup
       count at 1 for sessions with ~50K peak input.
    2) P's backup is a stale snapshot -- 49K of direct-to-D append work
       never flows through P. _commit_prefill_backup_residency
       (replay.py:1483) is only called from seed/reseed paths;
       direct-to-D path (replay.py:2719) never touches P-side state.
    3) When D evicts, old KV is freed directly (no D->P dump).
       session_aware_cache.release_session only calls
       kv_pool_allocator.free().

§2. End-to-end reseed timeline (t=0 to t=4550ms) with code citations
    showing exactly where each component sits. P-side re-prefill =
    1.5-3s, mooncake transfer = 1.5-4s, both contributing 50/50 to
    total reseed cost.

§3. Table of "looks like D->P but isn't" code locations -- every
    candidate found during forensic search ruled out with line citations.

§4. Specification of what D->P incremental sync would require:
    mooncake bidirectional roles (~400 LOC), D-side append commit hook
    (easy), P-side radix tree multi-producer extension (the real blocker),
    agentic-pd-hybrid replay.py hooks. Estimated 1-2 weeks engineering.

§5. Confirmation via `git ls-remote origin --refs` that author has NOT
    secretly implemented D->P on another branch -- only main + this
    working branch exist on the server.

§6. Roadmap for the upcoming feat/d-to-p-sync branch.

Appendices: code position crosswalk, related commits, paper section
suggestions.

This document is referenced by V2_DEEP_ANALYSIS_ZH §4.2 and by
KVC_ROUTER_ALGORITHM §9 Open Question 4.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

.gitignore

@@ -12,6 +12,7 @@ src/*.egg-info
 .deps/
 outputs/
 
-# Local heavyweight checkouts and generated experiment artifacts
-third_party/
+# Vendored dependencies. Track only the maintained SGLang fork/snapshot.
+third_party/*
+!third_party/sglang/
 *.log

AGENTS.md

@@ -0,0 +1,94 @@
+# AGENTS.md
+
+## Environment
+
+Use `uv` to manage all python environment. `uv add` to manage deps so that we can `uv sync` to get exactly same runnable environment each time.
+
+## Goal
+
+Build a minimal prototype on top of **SGLang xPyD** to test whether **session-aware / KV-cache-aware P/D routing** can improve **end-to-end latency** for agentic coding workloads.
+
+Current setup:
+- SGLang: `v0.5.10`
+- Model: `Qwen3-Coder-30B-A3B-Instruct` (`~/models/Qwen/Qwen3-Coder-30B-A3B-Instruct`)
+- xPyD runs on this single 8-GPU node, so the current constraint is **$x + y \le 8$**
+- Even in local experiments, the implementation should preserve the **P -> D RDMA-style data path** semantics as much as possible; local runs should treat this as a loopback-based stand-in rather than collapsing P/D into a special in-process shortcut
+- Traces:
+  - Ali coding agent (`~/ali-trace/trace-qwen3-coder-formatted/041715-041717.jsonl`)
+
+---
+
+## MVP Scope
+
+We only do the following:
+
+1. Run **SGLang xPyD** correctly on one machine
+2. Add a **baseline router**
+   - `turn1`: default routing
+   - `turn2+`: prefer previous `D` node for the same session
+3. Add a **KV-cache-aware routing** policy
+4. Replay traces and compare policies with the same evaluation pipeline
+
+Out of scope for now:
+- autoscaling
+- fault tolerance
+- large-scale cluster scheduler
+- production hardening
+- general multi-tenant serving
+
+---
+
+## What matters
+
+Primary metric:
+- **E2E latency**
+
+Secondary metrics:
+- TTFT
+- TPOT
+- KV transfer volume
+- cache hit / reuse
+- re-prefill count
+- per-node load
+
+Do not optimize TTFT alone if E2E does not improve.
+
+---
+
+## Development Order
+
+Implement in this order:
+
+1. **Bring up xPyD**
+2. **Add trace replay + metrics logging**
+3. **Implement sticky-to-D baseline**
+4. **Implement KV-cache-aware routing**
+5. **Analyze gains and failure cases**
+
+Do not skip step 2.
+
+---
+
+## Core Rules
+
+### 1. Keep policy separate from mechanism
+- mechanism = how requests / KV / xPyD work
+- policy = how we choose `P` and `D`
+
+Do not mix them unless necessary.
+
+### 2. Prefer simple, debuggable logic
+Start with simple heuristics before complex scoring.
+
+### 3. Log everything needed to explain results
+Each request should log:
+- request id
+- session id
+- turn id
+- assigned P node
+- assigned D node
+- latency
+- whether reuse was expected / observed
+
+### 4. Small interfaces only
+Avoid over-abstraction.

README.md

@@ -0,0 +1,101 @@
+# Agentic PD Hybrid
+
+这个项目是在 SGLang xPyD 上做一个最小实验框架，用来判断：
+
+**面向 agentic coding workload 的 session-aware / KV-cache-aware P/D routing，能不能降低端到端延迟。**
+
+更完整但仍然简洁的说明见 [docs/PROJECT_OVERVIEW.md](docs/PROJECT_OVERVIEW.md)。
+
+## 当前做了什么
+
+- 启动单机 SGLang P/D 栈。
+- 回放 Ali coding agent trace，并记录 request-level metrics。
+- 支持 `default`、`sticky`、`kv-aware` 路由策略。
+- 支持 `pd-disaggregation`、`kvcache-centric`、`pd-colo` 对比。
+- 支持小 append、多轮 session 的 micro-benchmark trace。
+- 维护了基于 SGLang `v0.5.10` 的本地 patch，放在 `third_party/sglang`。
+
+## 环境
+
+统一使用 `uv`：
+
+```bash
+uv sync
+```
+
+默认模型路径：
+
+```text
+~/models/Qwen/Qwen3-Coder-30B-A3B-Instruct
+```
+
+当前主要测试环境是单机 8 GPU，约束是 `prefill + decode <= 8`。
+
+## 常用命令
+
+生成小 append trace：
+
+```bash
+uv run agentic-pd-hybrid make-small-append-trace \
+  --output outputs/smoke-hotcap-30k-1k-256.jsonl \
+  --session-count 4 \
+  --turns-per-session 3 \
+  --initial-input-length 30000 \
+  --append-input-length 1000 \
+  --output-length 256
+```
+
+跑 live benchmark：
+
+```bash
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/micro-serveable-varturn-30k-1k-256-20260424T0756Z.jsonl \
+  --output-root outputs/live-serveable-varturn-30k-1k-256-hotcap \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --kvcache-admission-mode worker \
+  --prefill-workers 1 \
+  --decode-workers 1 \
+  --prefill-gpu-ids 0 \
+  --decode-gpu-ids 1 \
+  --transfer-backend mooncake \
+  --target-duration-s 2000 \
+  --session-sample-rate 1.0 \
+  --min-turns 2 \
+  --time-scale 1 \
+  --concurrency-limit 1000
+```
+
+只回放并写 metrics：
+
+```bash
+uv run agentic-pd-hybrid replay \
+  --trace path/to/trace.jsonl \
+  --policy kv-aware \
+  --mechanism pd-disaggregation \
+  --router-url http://127.0.0.1:8000 \
+  --output outputs/replay.jsonl
+```
+
+## 输出
+
+每次 replay/benchmark 会写：
+
+- request metrics：`request-metrics.jsonl`
+- 汇总结果：`request-metrics.jsonl.summary.json`
+
+重点看：
+
+- E2E latency
+- TTFT / TPOT
+- execution mode
+- cached tokens
+- KV transfer blocks
+- error
+
+## 维护约定
+
+- 项目代码改动：`feat:` / `fix:` / `docs:`。
+- SGLang 改动：`feat(sglang): ...` / `fix(sglang): ...`。
+- `third_party/sglang` 的基线是 clean SGLang `v0.5.10` snapshot。
+- 不提交 `outputs/`、日志、`__pycache__`、虚拟环境。

docs/AGENTIC_FIT_ANALYSIS_ZH.md

@@ -0,0 +1,434 @@
+# Agentic 场景下的结构性设计缺陷分析
+
+**日期**：2026-05-06
+**对照数据**：`outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/exp2_2p6d_run1_*`（KVC kv-aware Option D，2P6D，4449 reqs / 52 sessions）+ `outputs/qwen3-30b-tp1-exps/exp1_8way_dp_cache_aware_summary.json`（同 trace 8-way DP cache-aware baseline）。
+**模型**：Qwen3-30B-A3B（TP1），单机 8×H100 80GB。
+**研究问题**：把 SWE trace 视为"真实 agentic"的代表，KVC 机制相对 vanilla DP 系统性输在哪里——除了"D 容量 4.6× 过载"之外的结构性原因。
+
+> 本文是对 `docs/KVC_DEBUG_JOURNEY_V1_TO_V5.md` 与 `docs/V5_PROFILE_INVESTIGATION_ZH.md` 的补充：版本演进与瓶颈定位之外，从设计层看哪些假设和真实 agentic workload 不匹配。
+
+---
+
+## TL;DR
+
+按重要性排序的结构性缺陷：
+
+| # | 缺陷 | 数据 | 修复方向 | 工程量 |
+|---|---|---|---|---|
+| 1 | **KvAwarePolicy 不感知 D 容量；session 永久 pin 到首次落点 D** | session 平均访问的不同 D 数 = **1.00**；direct-to-D 命中率呈极端双峰（15 session 0-20%、14 session 80-100%） | score 函数加 capacity-aware 项；允许跨 D session 迁移 | 中 |
+| 2 | **D 端 LRU 只能 evict idle session，hot session 永远踢不掉** | D 跑全程仅 9-43 次 trim 事件 vs 80-150 次 transfer 错误；token_usage 顶到 1.00 | 加 score-based eviction（按访问频率/最近性多层） | 中 |
+| 3 | **没有 D→Router→Replay 的 backpressure 通道** | concurrency 一路 32 不降；D 失败时 replay 无感 | admission 响应加 `recommended_pause_ms`；replay 端按它降并发 | 小 |
+| 4 | **Admission HTTP round-trip 与 scheduler 主循环耦合** | v5+profile 仅加 1Hz polling 就让 errors 从 9 涨到 415 | 拆成 lock-free `/probe` + 进 scheduler 队列的 `/commit_evict` | 中 |
+| 5 | **P-side round-robin 不感知 D 健康** | prefill-0 出 367 KVTransferError，prefill-1 仅 4——但请求量近乎对半 | router 选 P 时考虑目标 D 健康度 | 中 |
+| 6 | **Replay 端 session footprint 估算膨胀 30×** | `_estimate_session_resident_tokens = input + output`，把 turn-50 的 80K 上下文当成"需要全新 80K 空间" | 改成"增量 token"估算 | 小 |
+| 7 | **time-scale=10 把测试条件人为推到失真区间** | inter-turn gap p50 从 2.5s 压到 0.25s——KVC 想利用的"自然 idle 窗口"被消除 | 跑一组 time-scale=1 baseline 验证 | 小（仅配置） |
+
+**最重要的对照事实**：同 trace、同硬件、同模型下 8-way DP cache-aware（无 PD 拆分、无 KVC、无 session 抽象）：
+
+| 指标 | 8-way DP CA | v5 KVC 2P6D |
+|---|---|---|
+| Errors | **0** | 372 (8.4%) |
+| Latency mean | **1.43s** | 3.50s |
+| Latency P50 | **0.65s** | 1.11s |
+| Latency P99 | **8.37s** | 20.37s |
+| TTFT mean | **0.12s** | 2.13s |
+| TTFT P90 | **0.26s** | 6.47s |
+| Per-worker 请求量分布 | 508–619（±10%） | 561–858（±26%） |
+
+**naive DP 在每一项都赢，包括 latency mean 的 145% 优势**。这定义了 KVC 在该 workload 下"必须超过"的基线。
+
+---
+
+## 1. Session 永久 pin 到 D + 容量盲选（最核心问题）
+
+### 1.1 现象
+
+每个 session 在整次运行中只访问 **1.00 个不同 D worker**（见上文数据）。结合 direct-to-D 命中率分布：
+
+```
+direct-to-D 命中率分桶（n=52 sessions）：
+  0-20%:  15 sessions ← 几乎每 turn 都失败回退到 P→D 全量传输
+  20-40%:  7
+  40-60%: 11
+  60-80%:  5
+  80-100%: 14 sessions ← 几乎每 turn 都走 direct-to-D 快路径
+```
+
+**几乎没有中间态**——这是典型的不公平资源分配信号。
+
+被饿死与被照顾的 session 在工作量上差异明显：
+- 饿死 session 平均 peak input：56,011 token
+- 顺利 session 平均 peak input：31,344 token（**1.8× 差距**）
+
+**大 session 倾向被饿死**——因为它们在容量已紧张的 D 上更容易触发 admission 拒。
+
+### 1.2 根因（代码级）
+
+`policies.py:166-172` `KvAwarePolicy.select`：
+
+```python
+score = (
+    overlap + sticky * self.sticky_bonus,    # 主项: 历史 KV overlap
+    sticky,                                   # 二级: 是否 last_decode_worker
+    inflight_penalty,                         # 三级: 当前 inflight 数（很小）
+    assignment_penalty,                       # 四级: 累计被分配数（更小）
+)
+```
+
+评分中**完全无 D 当前容量项**。Session X 第一次落到 D-2 时积累 hash_id 在 D-2 上；之后无论 D-2 多满，X 的 turn N+1 都会被打分到 D-2（因为 overlap 主导）。
+
+更糟的是 `RoutingState.decode_resident_blocks`（`policies.py:46`）从不缩减——即使 D 早 evict 了某些块，replay 仍认为它们在那。运行中期所有 D 的 overlap 集合都接近"trace 全部 hash_id"，policy 退化为纯 sticky。
+
+### 1.3 后果——具体到 session 的体验
+
+**饿死 session（如 session 50400，105 turns，0 次 direct-to-D）每 turn 流程**：
+
+1. policy 选 D（永远是同一个）
+2. admission 拒（D 容量已被占住）
+3. 走 fallback-session-cap → P 全量 prefill 50K-100K token
+4. mooncake 推 KV → D 仍无空间 → 32s timeout 或 KVTransferError
+5. 用户每 turn 体验 5-10s 延迟，反复出错
+
+**顺利 session（如 session 3840，118 turns，97% direct-to-D）每 turn 流程**：
+
+1. policy 选 D（永远是该 session 的初始 D）
+2. admission 通过（这个 session 一直占着这个 D 的 slot）
+3. direct-to-D：D 上 append-prefill 几百 token，零 P 介入、零 mooncake transfer
+4. TTFT 0.043s、E2E 0.495s
+
+**这不是"平均慢一点"，是结构性不公平**——SLO 视角下 P99 是被饿死那 15 session 的尾巴拉出来的。
+
+### 1.4 为什么 naive DP 反而赢
+
+8-way DP cache-aware 用纯 hash-based 路由，没有 session 抽象，没有 PD 拆分：
+
+- 每个请求按 prefix hash 路由到一个 worker → 同 session 的 turn 在 worker 上自然有 prefix 命中
+- 容量过载时 SGLang 自己的 radix cache + 调度器统一管 KV 池
+- 不存在 admission/fallback/reseed 路径
+- 不存在 mooncake transfer
+- per-worker 负载误差 ±10%（vs KVC ±26%），自动接近均衡
+
+**KVC 引入的 session affinity / KV 复用 / admission 三件套，在容量紧张时反而加剧了不均衡，没有任何一项能挽回 vs DP 的差距。**
+
+### 1.5 修复方向
+
+`KvAwarePolicy.select` 里加：
+
+```python
+# 当前 D 容量利用率（worker-mode admission 已经能查到）
+capacity_penalty = -worker_capacity_used_ratio[worker.worker_id]
+
+# 当多个 D 都有 overlap 时，按容量挑最空的；
+# 当某 D 容量 > 阈值时，禁止该 D 进入候选
+if worker_capacity_used_ratio[worker.worker_id] > HARD_CAP:
+    continue
+
+score = (
+    overlap_capped,                # overlap 但限幅，避免单个 D 永远赢
+    capacity_penalty,              # ← 新增
+    sticky,
+    inflight_penalty,
+)
+```
+
+更激进的修法：当一个 session 被某 D 反复拒 N 次后，主动 release 它在该 D 上的 session 状态，**允许下次 turn 走另一个 D**（代价是丢失已积累的 KV，但目前 fallback 路径本来也丢了）。
+
+---
+
+## 2. D 端 LRU eviction 跟不上压力
+
+### 2.1 数据
+
+每个 D 全程：
+
+| Worker | Trim 事件（主动 LRU） | KVTransferError + OOM | 峰值 token_usage |
+|---|---:|---:|---:|
+| decode-0 | 9 | 0 | 0.99 |
+| decode-1 | 43 | 12 (4 err + 8 oom) | 0.99 |
+| decode-2 | 16 | 459 (153 err + 306 oom) | 0.97 |
+| decode-3 | 37 | 87 (29 err + 58 oom) | 0.99 |
+| decode-4 | 28 | 270 (90 err + 180 oom) | **1.00** |
+| decode-5 | 30 | 279 (93 err + 186 oom) | **1.00** |
+
+**LRU 触发频率比错误次数低 5-15 倍。** D-4 / D-5 直接顶到 token_usage=1.00。
+
+### 2.2 根因
+
+`scheduler.py:2040` `evict_idle_streaming_sessions_lru` 的 idle 判定：
+
+```python
+# 只能 evict "所有 req 都 finished + streaming 模式" 的 session
+```
+
+但 SWE 高并发下每个 session 几乎一直有 inflight req（time-scale=10 又压缩了 inter-turn gap）。**hot session 永远不 idle，LRU 永远找不到东西可踢**。结果 D 一路开到 100% → 下一笔 transfer 来直接 OOM/timeout。
+
+### 2.3 修复方向
+
+引入分层 eviction：
+
+1. **Idle session 优先**（当前）
+2. **冷 session 次优**（最近 N 秒无访问，即使有 inflight，也可以 retract 那个 inflight 让位）
+3. **hot session 强制 retract**（在 hard cap 触发时）
+
+vanilla SGLang 已有 `disagg_decode_prealloc_queue.retracted_queue` 机制（看 `admit_direct_append` 引用），但**没有人主动触发 retract**——目前只有内部异常时才会进 retracted_queue。需要把 retract 提升为正常 admission 路径的一部分。
+
+---
+
+## 3. 没有 D→Replay 的 backpressure 通道
+
+### 3.1 名词解释
+
+**Backpressure（反压）** = 流式系统下游过载时把信号反向传给上游让它降速。例：TCP 滑动窗口、Kafka consumer lag、gRPC HTTP/2 flow control。
+
+### 3.2 当前状态
+
+- D 端 transfer queue 堆 → 32s 后 timeout → 抛 KVTransferError
+- error 抛回 P → P 抛给 router → router 抛给 replay → replay 走 fallback 路径
+- **整个链路上没有"D 过载，请慢点发"的信号**——concurrency 一直保持上限
+
+后果：D 一旦开始失败，会**持续失败**（因为 replay 没降速），直到 D 自己消化完积压。
+
+### 3.3 修复方向
+
+`admit_direct_append` 响应里加：
+
+```python
+{
+  "can_admit": ...,
+  "recommended_pause_ms": int,    # ← 新增：下次发同类请求前建议等多久
+  "queue_depth": int,             # ← 新增：D transfer queue 当前深度
+  ...
+}
+```
+
+replay 端在 admission 拒被拒时按 `recommended_pause_ms` 降并发或退避。**这是最便宜的一条改动**——不改协议、不改 SGLang 内部，只改两端代码。
+
+---
+
+## 4. Admission RPC 与 scheduler 耦合——结构 vs 工程的精确边界
+
+### 4.1 现象
+
+`docs/V5_PROFILE_INVESTIGATION_ZH.md` 报告：仅加 1Hz `/server_info` polling 就让 EXP2 errors 从 9 涨到 415。`/server_info` 在 scheduler 主循环里遍历 session slots 算 `is_idle`，1 Hz × 8 worker 就足以扰动调度。
+
+但实际负载下 admission RPC 频率远高于 1Hz：每个 turn 1 + reseed + direct-to-D 都调一次。concurrency=32 + 4449 reqs / ~2700s ≈ **每秒 16+ 次 admission RPC**。
+
+### 4.2 这是结构问题还是工程问题——精确拆解
+
+`admit_direct_append`（`scheduler.py:3581`）做两件事：
+
+```python
+# (a) 读池子状态——轻
+available_tokens = self.token_to_kv_pool_allocator.available_size()
+
+# (b) 触发 LRU 扫描——重，且必须修改池子状态
+trim_result = self.maybe_trim_decode_session_cache(...)
+```
+
+| 部分 | 性质 | 是否能靠工程化解决 |
+|---|---|---|
+| (a) 读池子状态 | 几个原子读 | **完全可工程化**——做成 lock-free shared-memory snapshot 即可 |
+| (b) LRU eviction | 修改 GPU 池子，必须独占 | **结构性的**——Python GIL + 共享 GPU 池子无法并发修改 |
+
+**关键观察**：实际负载里 (b) 是少数路径——大部分 admission 只需要"看一下够不够"，不需要立即 evict。
+
+### 4.3 工程化修复方案
+
+把 admission API 拆成两个端点：
+
+```
+POST /session_cache/probe          ← 90% 流量
+  - 只读 lock-free snapshot
+  - 返回 (can_admit_estimate, available_tokens, queue_depth)
+  - 不进 scheduler 队列
+
+POST /session_cache/commit_evict   ← 10% 流量
+  - probe 不够时才调
+  - 进 scheduler 队列，做实际 LRU
+  - 保留当前 admit_direct_append 语义
+```
+
+snapshot 由 scheduler 在每个 step 末尾写到一段 mmap 共享内存（atomic publish）；replay 端 mmap 读，零 syscall 零序列化。一秒内能撑数千次 probe。
+
+### 4.4 关于"协程/多线程/多进程/换语言"
+
+| 工具 | 对本问题的实际效果 |
+|---|---|
+| asyncio 协程 | SGLang 已用，对 scheduler 主循环本身无帮助 |
+| Python 多线程 | GIL 拦着，且 GPU 池子状态只能 scheduler 进程改 |
+| 多进程 | scheduler 已是独立进程；问题是它**自己的 step 循环**串行了 admission 与 decode |
+| orjson / uvloop | 网络/JSON 加速 5-10×，但 LRU 遍历不在那条热路径 |
+| Rust/C++ 重写 scheduler | 把 LRU 遍历提速 5-10×，但**结构性共享问题仍在** |
+
+**正确的工程化解法是重设计 API（拆 probe / commit），不是单纯换更快的库或语言。**
+
+---
+
+## 5. P-side 路由不感知 D 健康
+
+### 5.1 数据
+
+```
+prefill-0:  367 KVTransferError, 361 "Decode instance could be dead"
+prefill-1:    4 KVTransferError, 0  "Decode instance could be dead"
+
+请求量对比:
+  prefill-0: 2225 requests
+  prefill-1: 2224 requests   ← 几乎对半
+```
+
+**两 P 请求量完全均衡，错误率差 92×**。日志里 prefill-0 的错误反复指向某个特定 D（`10.45.80.47:XXXXX`）——它跟某个 hot D 形成了"死亡链路"。
+
+### 5.2 根因
+
+`pd_router.py:43-49` 的 P 选择是裸 round-robin：
+
+```python
+prefill_url, bootstrap_port = self.config.prefill_urls[
+    self.prefill_cursor % len(self.config.prefill_urls)
+]
+```
+
+不知道 D 是否健康，不会避开"正在和 D-X 死磕"的 P。
+
+### 5.3 修复方向
+
+router 选 P 时考虑 (P 当前 inflight transfer 数, 目标 D 健康度) 联合得分。健康度可以用 §3 提的 `queue_depth` 字段。
+
+---
+
+## 6. Replay 端 session footprint 估算膨胀 30×
+
+### 6.1 代码
+
+`replay.py:898-899`：
+
+```python
+def _estimate_session_resident_tokens(request: TraceRequest) -> int:
+    return request.input_length + request.output_length
+```
+
+被用于 `_decode_session_soft_cap`（`replay.py:1051`）和 `_should_admit_new_decode_session`。
+
+### 6.2 问题
+
+对一个已经在 D 上有 80K KV 的 turn 50：
+- 真实增量需求：input 新增几千 token + output 几百 token = ~3K
+- 估算返回值：80K + 1K = 81K（**膨胀 ~27×**）
+
+后果：router-mode admission 系统性误判——本来能 admit 的 session 被 replay 自己拒掉。v5 worker-mode 让 D 自己看真实容量部分修了这个，**但 KvAwarePolicy 选 D 时仍用这个膨胀估算**——选 D 仍然是错的。
+
+### 6.3 修复
+
+```python
+def _estimate_session_resident_tokens(request: TraceRequest) -> int:
+    if request.turn_id == 1:
+        return request.input_length + request.output_length
+    # turn 2+: only the increment matters for additional reservation
+    return max(0, request.input_length - request.cached_tokens) + request.output_length
+```
+
+---
+
+## 7. time-scale=10 测量失真
+
+### 7.1 它是什么
+
+`replay.py` 把原始 trace 每个请求的 `timestamp` 字段做 `t / time_scale` 缩放后再按这个时间发。
+
+- 原始 trace 跨度 ~6000s（≈100 分钟）
+- time-scale=10 → 实际 replay 跨度 ~600s（≈10 分钟）
+
+### 7.2 为什么这么设计
+
+**纯粹为了节省测试时间**——单次 1× 跑 100 分钟，sweep 5 版 × 3 重复 = 25h GPU 时间；10× 只要 2.5h。
+
+### 7.3 它扭曲了什么
+
+| 维度 | 原始 trace | replay (time-scale=10) |
+|---|---|---|
+| inter-turn gap p10 | 1.6s | 0.16s |
+| inter-turn gap p50 | 2.5s | 0.25s |
+| inter-turn gap p90 | 7.8s | 0.78s |
+| inter-turn gap max | 261s | 26s |
+
+真实 agentic 用户/agent 在每个 turn 之间停 2-8 秒（思考、打字、tool call）。**这些间隙正好是 KVC 想利用的"自然 idle 窗口"**——session 短暂 idle 时 LRU 可以 evict、其他 session 可以 admit。
+
+time-scale=10 把这些窗口压到 0.2-0.8s，**人为消除了 KVC 的设计前提条件**。
+
+### 7.4 严重的实验有效性威胁
+
+所有 v3-v6 数据基于 time-scale=10。这意味着前面所有"KVC 在 SWE 上输给 baseline"的结论都带着这个失真。**真实部署里 inter-turn gap 是 2.5s 的话，KVC 可能根本不会撞到当前看到的容量瓶颈**——D 有时间在 turn 之间释放/重排。
+
+**应该单独跑一组 time-scale=1 的 baseline 对比**，才能判断 KVC 输给 DP 是因为机制本身不行，还是因为 benchmark 把它推到了不该工作的区间。这是这个项目目前**最重要但还没做**的验证。
+
+---
+
+## 8. 应用层抽象不需要在引擎层引入（撤回）
+
+之前草稿里提过"框架不支持 speculative 多分支、嵌套 sub-agent、tool call 中断"——这是过度抽象。**应用层模式都可以由 timestamp + 独立 session_id 隐式表达**：
+
+| 应用层模式 | 表现在 trace 里 | 推理引擎需要做什么 |
+|---|---|---|
+| Tool call 异步返回 | turn N 与 N+1 之间 timestamp gap 很大 | 啥都不用，按时间发请求即可 |
+| 嵌套 sub-agent | 父 session timestamp 突然停顿；sub-agent 是独立 session_id | 把它们当成两个独立 session 即可（KV 也无需共享） |
+| Speculative N 分支 | N 个独立 session_id 同时发 | 用 radix prefix cache 自然命中前缀；不需要任何额外抽象 |
+
+**这条不构成结构性缺陷。** 已从结论中移除。
+
+---
+
+## 9. 行动项（按 ROI 排序）
+
+### 优先级 P0（修了显著改善饿死/不公平）
+
+1. **[§1] KvAwarePolicy 加 capacity-aware penalty + 允许 session 跨 D 迁移** — 工程量中、收益最大
+2. **[§2] D 端引入分层 eviction（冷 session、hot retract）** — 工程量中、收益大
+3. **[§7] 跑一组 time-scale=1 baseline** — 工程量小（仅配置），但**不做这条所有结论都不可信**
+
+### 优先级 P1（修了把工程稳定性补齐）
+
+4. **[§3] D→Replay backpressure 通道**（admission 响应加 pause hint） — 工程量小
+5. **[§4] 拆 admission 为 probe + commit_evict** — 工程量中
+6. **[§6] 修 `_estimate_session_resident_tokens` 用增量** — 工程量小
+
+### 优先级 P2（等 P0 数据后再决定）
+
+7. **[§5] P-side 选 P 时考虑 D 健康** — 工程量中
+
+---
+
+## 10. 局限与未验证假设
+
+1. **N=1**：所有数据来自单次 run（v6 P0 已证 EXP2 errors 在 9-912 间漂移，single-run variance 巨大）。本文所有数字都应理解为"代表性观察"而非"统计显著结论"。
+2. **time-scale=10 失真**（§7）：所有"KVC 输给 DP"的程度可能是被 benchmark 放大的。这是最大的不确定性。
+3. **8DP 对比的硬件优势**：DP 是 8 个 worker 全部跑 prefill+decode；KVC 是 2P+6D，只有 6 个能解码。理论上 8 worker 对 6 worker 自带 1.33× 解码并发优势。本文未折算这部分——但 8DP 优势远大于 1.33×（latency mean 145% 优势），所以核心结论（KVC 在该 workload 下系统性输）不受此影响。
+4. **mooncake TCP loopback**：所有 transfer 错误是单机 TCP 模拟下的产物。生产环境 RDMA 下错误率分布可能完全不同。
+5. **KvAwarePolicy 的 stale `decode_resident_blocks`**（§1.2 末尾）现象有数据观察支撑（运行中期 overlap 失去判别力），但**没有系统性测过"清掉 stale 状态会怎样"**。
+6. **P-side 错误集中在 prefill-0**（§5.1）的因果链是推测——可能也是"prefill-0 早启动 + race"的偶然结果。N>1 数据未验证。
+
+---
+
+## 附录 A：数据产物索引
+
+```
+outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/
+├── exp2_2p6d_run1_metrics.jsonl    ← 本文主数据源
+├── exp2_2p6d_run1_summary.json
+├── exp2_2p6d_run2_*  (errors=912, single-run variance 证据)
+├── exp2_2p6d_run3_*  (errors=396)
+└── kvcache-centric-*-20260429T142429Z/logs/
+    ├── decode-{0..5}.log           ← §2.1 LRU vs error 计数
+    └── prefill-{0,1}.log           ← §5.1 P 错误分布
+
+outputs/qwen3-30b-tp1-exps/
+├── exp1_8way_dp_cache_aware_summary.json   ← 对照 baseline
+└── RESULTS_SUMMARY.md
+```
+
+## 附录 B：相关文档
+
+- `docs/PROJECT_OVERVIEW.md` — 项目目标与已实现功能
+- `docs/KVC_DEBUG_JOURNEY_V1_TO_V5.md` — v1→v5 版本演进
+- `docs/V5_PROFILE_INVESTIGATION_ZH.md` — v5+profile 调查（已 critic 修订）
+- `docs/SWEBENCH_EXPERIMENT_RESULTS.md` — Qwen3.5-35B-A3B SWE 实验

docs/KVCACHE_CENTRIC_PROGRESS_ZH.md

@@ -0,0 +1,263 @@
+# KV-cache centric P/D routing 当前进展
+
+本文记录当前原型在 SGLang xPyD 上围绕 session-aware / KV-cache-aware P/D routing 的实现、实验结果和阶段性结论。实验日期为 2026-04-24 至 2026-04-25。
+
+## 目标和核心假设
+
+目标是在单机 8 GPU xPyD 环境中验证：针对 agentic coding workload，session-aware / KV-cache-aware P/D routing 是否能提升端到端延迟。
+
+当前重点假设：
+
+1. 在 PD-disaggregation 下，P 节点和 D 节点可能同时保留同一个 session 的 prefix KV，形成 P/D duplicate。
+2. 如果预测某个 session 后续会 direct-to-D，那么 P 侧在 radix/prefix cache eviction 时可以优先淘汰这部分 prefix cache。
+3. 这样可以给 P 节点释放 cache 空间，提高 P 侧 prefix cache reuse。
+4. 如果 P 侧 reuse 提升后 D 侧开始成为瓶颈，可以通过增加 xPyD 中 D 的配比，也就是增加 y，缓解 decode 侧压力。
+
+实验结果表明：第 2、3 点在 P cache 高压 workload 下成立；第 4 点只部分成立，因为瓶颈会从 decode prealloc 转移到 P->D transfer/bootstrap pipeline。
+
+## 已实现机制
+
+### 1. Trace profile 和 paired comparison
+
+新增 `agentic_pd_hybrid profile` 子命令和 `src/agentic_pd_hybrid/profile.py`。
+
+能力：
+
+- 统计 trace 的 session 数、turn2+ 数、append tokens、overlap ratio、direct-to-D eligible turn。
+- 对比 baseline/candidate metrics，输出 paired E2E latency delta。
+- 用于解释 micro-benchmark 与 Ali filtered workload 的差异。
+
+### 2. P 侧 priority eviction 支持
+
+修改 SGLang server args，允许：
+
+```bash
+--radix-eviction-policy priority
+```
+
+router/replay 支持内部字段：
+
+- `smg_prefill_priority`
+- `smg_decode_priority`
+
+router 会将内部字段剥离，只把标准 `priority` 分别传给 P/D backend。这样可以让 direct-to-D predicted session 在 P 侧使用更低优先级，例如 `-100`，普通请求使用 `100`。
+
+### 3. Kvcache seed/direct admission 控制
+
+新增多种 seed/reseed 过滤：
+
+- `kvcache_seed_max_resident_tokens`
+- `kvcache_seed_max_output_tokens`
+- `kvcache_seed_min_turn_id`
+- `kvcache_seed_only_multiturn_sessions`
+- `kvcache_seed_max_inflight_decode`
+
+新增 P streaming session backup 策略：
+
+- `release-after-transfer`：P->D transfer 后释放 P 侧 session backup。
+- `capacity-backup`：容量允许时保留 P 侧 backup。
+
+当前主实验使用 `release-after-transfer`。
+
+### 4. 稳定性修复
+
+之前 worker-admission 实验会在 replay 尾部卡住，不写出 metrics。原因是 benchmark 的长 `timeout_s=3600` 同时用于：
+
+- SGLang stack 启动等待
+- replay client 单请求
+- router 到 P/D backend 的单请求
+
+修复后新增：
+
+- `BenchmarkConfig.request_timeout_s`
+- CLI `benchmark-live --request-timeout-s`
+- launch plan `router_request_timeout_s`
+
+当前做法：
+
+- `timeout_s=3600` 继续用于 SGLang 启动和整体 stack 等待。
+- `request_timeout_s=180` 用于 replay client 和 router 到 P/D backend 的单请求。
+- control-plane probe/open/close session 使用 2s timeout，fail closed。
+
+修复效果：worker-admission 从“尾部卡死不落盘”变为“卡住请求记录为 ReadTimeout，整轮实验完成并写 metrics”。
+
+## 关键实验结果
+
+### P cache pressure 下 priority eviction 是否提升 P 侧 reuse
+
+配置：2P2D，P `--max-total-tokens 90000`，micro workload 316 requests / 58 sessions。
+
+| 配置 | ok/total | mean E2E | p99 | request cached tokens | P log cached tokens | P new-token total |
+|---|---:|---:|---:|---:|---:|---:|
+| LRU | 314/316 | 28.171s | 43.409s | 8.204M | 7.783M | 2.236M |
+| Priority | 314/316 | 28.165s | 41.935s | 8.401M | 7.981M | 2.039M |
+
+结论：
+
+- 在 P cache eviction 高频触发时，priority eviction 确实提高 P 侧 prefix reuse。
+- cached tokens 增加约 197k，new prefill tokens 减少约 197k。
+- 但 mean E2E 基本不变，说明性能瓶颈转移到 D decode/transfer。
+
+### 增加 D 配比前的 D 侧瓶颈证据
+
+2P2D priority pressure 下：
+
+- decode `#queue-req`：max/mean/p90 = 0/0/0
+- decode token usage：max 0.98，mean 0.842，p90 0.96
+- decode `#transfer-req`：max 7，mean 4.61，p90 7
+- decode `#prealloc-req`：max 21，mean 11.6，p90 21
+
+解释：
+
+- D 侧不是普通 waiting queue 堆积。
+- 真正压力在 token usage、transfer queue 和 prealloc queue。
+
+### D scaling：router admission 旧结果
+
+| 配置 | ok/total | mean | p50 | p90 | p99 | error |
+|---|---:|---:|---:|---:|---:|---:|
+| 2P2D | 314/316 | 28.165s | 30.576s | 38.267s | 41.935s | 2 |
+| 2P3D | 290/316 | 29.915s | 31.428s | 40.856s | 45.964s | 26 |
+| 2P4D | 285/316 | 30.566s | 32.823s | 40.566s | 44.838s | 31 |
+
+结论：
+
+- 直接增加 D 不稳定。
+- 2P3D/2P4D 的错误主要来自 `kvcache-centric` seed/direct 路径。
+- 日志显示 decode 侧出现 `WaitingForInput` timeout 和 `KVTransferError`。
+
+### D scaling：worker admission + request timeout 修复后
+
+配置：P `--max-total-tokens 90000`，priority eviction，worker admission，`request_timeout_s=180`。
+
+| 配置 | ok/total | mean | p50 | p90 | p99 | error |
+|---|---:|---:|---:|---:|---:|---:|
+| 2P2D | 313/316 | 29.838s | 30.742s | 39.641s | 52.506s | 3 |
+| 2P3D | 299/316 | 29.349s | 30.569s | 42.161s | 46.113s | 17 |
+| 2P4D | 312/316 | 26.442s | 27.759s | 38.197s | 47.970s | 4 |
+
+对应 decode log 摘要：
+
+| 配置 | decode usage mean | decode transfer mean | decode prealloc mean |
+|---|---:|---:|---:|
+| 2P2D | 0.859 | 4.86 | 11.3 |
+| 2P3D | 0.877 | 5.70 | 6.92 |
+| 2P4D | 0.809 | 5.31 | 3.46 |
+
+结论：
+
+- 2P4D + worker admission + request timeout 是当前最好的 D scaling 配置。
+- 相比旧 2P4D，成功率从 285/316 提升到 312/316，mean 从 30.566s 降到 26.442s。
+- 但 p99 仍未稳定改善，tail 仍由 P->D transfer/bootstrap timeout 主导。
+- 2P3D 不稳定，错误 17 个，不适合作为当前推荐配置。
+
+### 与默认 PD-disaggregation 的同配置对比
+
+为了避免和旧 2P2D/no-pressure 基线混比，补跑了同一 workload、同一 2P4D、同一 P `--max-total-tokens 90000`、同一 `time_scale=50`、同一 `request_timeout_s=180` 的默认 PD-disaggregation 基线。
+
+| 策略 | ok/total | mean | p50 | p90 | p99 | cached total | direct-to-D |
+|---|---:|---:|---:|---:|---:|---:|---:|
+| 默认 PD-disaggregation 2P4D | 316/316 | 29.210s | 28.940s | 47.434s | 52.605s | 8.165M | 0 |
+| KVC 2P4D latency-best | 312/316 | 26.442s | 27.759s | 38.197s | 47.970s | 7.882M | 11 |
+| KVC 2P4D seed-min2 | 316/316 | 26.729s | 25.840s | 43.589s | 50.426s | 8.337M | 3 |
+
+结论：
+
+- 如果只看成功请求延迟，KVC 2P4D latency-best 相比默认 PD：
+  - mean 改善约 9.5%；
+  - p50 改善约 4.1%；
+  - p90 改善约 19.5%；
+  - p99 改善约 8.8%。
+- 但 latency-best 有 4 个 `ReadTimeout`，全部来自 turn1 大 seed 的 P->D transfer/bootstrap timeout。
+- `kvcache_seed_min_turn_id=2` 可以消除这些错误，达到 316/316 成功，同时 mean/p50/p90/p99 仍然优于默认 PD。
+- 因此当前推荐分成两档：
+  - 追求最低 mean/p90：使用 KVC 2P4D latency-best。
+  - 追求稳定性：使用 KVC 2P4D seed-min2。
+
+### 进一步优化尝试
+
+在当前最佳配置基础上尝试了三个优化方向：
+
+| 策略 | ok/total | mean | p50 | p90 | p99 | 结论 |
+|---|---:|---:|---:|---:|---:|---|
+| KVC 2P4D latency-best | 312/316 | 26.442s | 27.759s | 38.197s | 47.970s | 延迟最优，但有 4 个 turn1 seed timeout |
+| transfer-cap4 | 304/316 | 27.961s | 29.785s | 38.730s | 45.192s | 不可取，实时 transfer queue snapshot 滞后，错误更多 |
+| disable-failed-session | 285/316 | 29.179s | 30.878s | 42.794s | 46.411s | 不可取，失败 session 降级会放大后端异常状态 |
+| seed-min2 | 316/316 | 26.729s | 25.840s | 43.589s | 50.426s | 稳定性最优 |
+| inflight0 | 316/316 | 29.497s | 31.186s | 43.498s | 48.022s | 太保守，几乎关闭 KVC 收益 |
+
+从错误明细看，latency-best 的 4 个错误全部是 turn1 seed timeout，每个都是约 1250 KV blocks 的大 seed。`seed-min2` 跳过 turn1 seed 后，错误全部消失，说明当前主要稳定性问题是启动阶段 seed 风暴，而不是后续 direct append 本身。
+
+transfer-cap4 没有效果，说明仅依赖 D worker 的实时 `decode_transfer_queue_reqs` 不够；并发请求可能同时读取到尚可的 snapshot，然后一起进入 P->D transfer，导致 backlog 在 admission 后形成。
+
+## Ali filtered 当前状态
+
+Ali filtered small-append trace：
+
+- 81 requests
+- 28 sessions
+- 53 turn2+
+- max input 18901
+- max output 1925
+- span 5414s
+
+PD baseline：
+
+- ok 81/81
+- mean 9.072s
+- p50 7.086s
+- p90 21.761s
+- p99 26.813s
+
+Kvcache-centric 在 Ali filtered 上曾出现 58-67 个 router 200 后挂住、无 metrics 的问题。当前 request timeout 和 control-plane timeout 修复后，应重新跑 Ali filtered；在未重跑前，不把 Ali filtered 纳入最终性能结论。
+
+## 当前结论
+
+1. kvcache centric 可以提高 KV reuse，但需要满足 workload 条件：
+   - session 有多 turn；
+   - turn2+ append 较小；
+   - prefix overlap 高；
+   - P 侧 cache 有 eviction pressure；
+   - D 侧 seed/direct admission 不把 transfer pipeline 打爆。
+
+2. 不适合的 workload：
+   - 单 turn 或 session 间隔过长；
+   - turn2+ append 很大，direct-to-D 不能省掉多少 prefill；
+   - prefix overlap 低；
+   - P cache 没有 eviction pressure；
+   - D transfer/prealloc 已经高压。
+
+3. 用户关于 P/D duplicate 的假设部分成立：
+   - 如果 D session 已经 resident，P 侧对应 streaming session backup 可以视为 duplicate。
+   - `release-after-transfer` 可以避免长期保留 P/D 两份 session KV。
+   - priority eviction 进一步让 P 在必须 eviction 时优先淘汰 direct-to-D predicted session prefix。
+
+4. 但当前机制还没有完全解决性能问题：
+   - P 侧 reuse 提升后，E2E 不一定改善。
+   - 主要原因是 D 侧 transfer/bootstrap pipeline 成为瓶颈。
+   - 增加 D 可以降低 prealloc，但不能自动降低 transfer backlog。
+   - 当前同配置下 KVC 已经可以优于默认 PD，但必须在 latency-best 和 stable 两种策略之间取舍。
+
+## 下一步优化方向
+
+1. transfer-aware admission：
+   - seed/direct 不只看 D token capacity，也要看 `decode_transfer_queue_reqs`、`decode_prealloc_queue_reqs`、`decode_retracted_queue_reqs`。
+   - 当 transfer queue 高时，应该主动走 PD fallback。
+   - 当前实验显示，单纯使用实时 transfer queue threshold 不够，需要配合本地 reservation。
+
+2. per-D transfer budget：
+   - 对每个 D 设置 seed/reseed 并发上限。
+   - 不能只按 session residency 或 token headroom 判断。
+   - 特别要限制 turn1 大 seed 的并发，避免启动阶段 seed 风暴。
+
+3. P/D ratio 联合调度：
+   - 2P4D 当前最好，但 P queue 也随 D 增加而上升。
+   - 后续需要测试 3P3D、3P4D、4P4D 等组合，确认 P transfer source 是否成为瓶颈。
+
+4. Ali filtered 重跑：
+   - 使用 request timeout 修复后的版本重新跑 Ali filtered。
+   - 如果仍然没有收益，需要按 session gap、append size、overlap ratio 分桶分析。
+
+5. 更严格的成功率指标：
+   - 当前不能只看成功请求的 mean/p90。
+   - 必须同时报告 ok/total、timeout/error 类型和 tail latency。

docs/KVC_DEBUG_JOURNEY_V1_TO_V5.md

@@ -0,0 +1,367 @@
+# KVC 实验踩坑记录与代码 Bug 分析（v1 → v5）
+
+记录从 v1 到 v5 KVC 实验的踩坑过程、错误诊断、以及最终定位的代码 bug。
+模型: Qwen3-30B-A3B (TP1)，硬件: 单节点 8×H100 80GB。
+Trace: `qwen35-swebench-50sess.jsonl`（4449 请求，52 sessions）。
+
+## TL;DR
+
+| 版本 | 关键变化 | 截断率 | direct-to-D 占比 | P50 | 主要瓶颈 |
+|------|----------|:---:|:---:|:---:|----------|
+| v1 (smoke / 早期) | mechanism 跑通 | - | - | - | - |
+| v2 | KVC + `--policy default` | **56.8% / 61.4%** | <0.1% | 0.08s* | Routing 错位（默认策略） |
+| v3 | KVC + `--policy kv-aware` | **0.9%** | 30-42% | 1.5-1.8s | session-cap fallback (52-65%) |
+| v4 | v3 + soft_cap 4→16 | 1.0% | 54-58% | 1.08 / 0.84s | session-cap fb 仍 35%、9-10% mooncake errors |
+| v5 | Option D：worker-mode 驱动 seed/reseed | 0.9% | 41-45% | 1.59 / 1.31s | D KV pool 真容量不足 → fallback 反而 ↑ 至 46-51% |
+
+`*` v2 的 P50 是假数字——超过半数请求只生成 1 个 token 就被 abort。
+
+## v2 踩坑：Default policy 与 KVC 机制根本不兼容
+
+### 表象
+
+`scripts/sweep_tp1_v2_fixed.sh` 跑出来：
+- Exp1（8-way DP，baseline）：4449/4449 成功，P50=0.65s，error=0
+- Exp2（1P7D KVC）：**2524 truncated (56.8%)**，18 errors，P50=0.08s* (假)
+- Exp3（2P6D KVC）：**2733 truncated (61.4%)**，17 errors，P50=0.08s* (假)
+
+每个截断请求 `actual_output_tokens=1`，`finish_reason="abort: session id X does not exist"`。
+
+### 错误的早期诊断
+
+之前 `RESULTS_SUMMARY.md` 把锅扣在 SGLang 的 `--disaggregation-decode-allow-local-prefill` flag 上，认为是 D worker 在有 `bootstrap_room` 时仍然做了 local prefill。这个诊断**完全错误**——查 `scheduler.py:1975-1980` 的 `_should_allow_local_prefill_on_decode`：
+
+```python
+def _should_allow_local_prefill_on_decode(self, req: Req) -> bool:
+    return (
+        self.disaggregation_mode == DisaggregationMode.DECODE
+        and self.server_args.disaggregation_decode_allow_local_prefill
+        and req.bootstrap_room is None  # ← 有 bootstrap_room 不会走 local prefill
+    )
+```
+
+KVC reseed 路径的请求都带 `bootstrap_room`，根本不会触发 local prefill。
+
+### 实际根因：Replay 与 PD Router 的 round-robin 错位
+
+实验脚本里 KVC 用 `--policy default`，而 baseline 用 `--policy kv-aware`。
+看 `benchmark.py:287-300` 这两者的差别巨大：
+
+```python
+def _decode_policy_for(policy_name: str) -> str:
+    if policy_name == "sticky":      return "manual"
+    if policy_name == "kv-aware":    return "consistent_hashing"
+    return "round_robin"  # default
+
+def _header_mode_for(policy_name: str) -> str:
+    if policy_name == "sticky":      return "routing-key"
+    if policy_name == "kv-aware":    return "target-worker"
+    return "none"  # default
+```
+
+`default` policy + KVC 机制下：
+1. Replay policy（`policies.py:DefaultPolicy`）round-robin 选一个 D，比如 D-3
+2. Replay 在 D-3 上 `open_session(session_id=X)`（`replay.py:1722-1731`）
+3. Replay 通过 PD Router 发请求（带 `session_params`），但 `header_mode=none`，**不发任何 routing header**
+4. PD Router (`pd_router.py:_select_decode_index`) 看到 `decode_policy=round_robin`，用**自己独立的计数器**round-robin，发到了 D-5
+5. D-5 的 scheduler 看到 `session_params` 里有 session_id，但自己的 `session_controller` 里没这个 session（session 在 D-3 上）→ abort with `"Invalid request: session id X does not exist"` (`scheduler.py:1824-1836`)
+
+两个独立的 round-robin 计数器只要一次错位（任何并发或 direct-to-D 绕过 router 的请求都会引起）就永远对不上。
+
+### 为什么 turn 0 不出问题？
+
+Turn 0 走 `_invoke_plain_router`（`replay.py:1894`），不带 `session_params`，作为普通 PD disagg 请求处理，发到任何 D 都行。Turn 1+ 才开始走带 session_params 的 KVC 路径，撞上路由错位。
+
+### 数据特征验证（per-session pattern）
+
+```
+session 11360 (58 turns): pattern = .TTTTT.TTTTTTT.TTTTTT...   ← turn 0 OK，1+ 全 T
+session 18720 (87 turns): pattern = .TTTTTTTTTTTTTTTTTT...
+```
+
+每个 D worker 收到了全部 52 个 session 的请求（理想情况下应该是 ~7-8 个/D，因为 round-robin 把 session 完全打散）。
+
+### 修复
+
+唯一正确的修复是把 KVC 的 policy 从 `default` 改成 `kv-aware`：
+
+```diff
+- --policy default
++ --policy kv-aware
+```
+
+`KvAwarePolicy` (`policies.py:146-187`) 做两件事：
+1. 用 `_overlap_blocks` + `sticky_bonus` 给每个 D 打分，session 自然粘在同一个 D（**session 亲和性**）
+2. `header_mode=target-worker`，发 `x-smg-target-worker` header
+3. PD Router 用 `consistent_hashing` 模式，看到 header 就直接用，不再 round-robin
+
+## v3 改 kv-aware policy 后：路由对了，但新瓶颈出现
+
+`scripts/sweep_tp1_v3_kvaware.sh` 把所有 KVC 实验改成 `--policy kv-aware`，结果：
+
+| 指标 | v2 1P7D (default) | **v3 1P7D (kv-aware)** | v3 2P6D | 8-way DP baseline |
+|------|:---:|:---:|:---:|:---:|
+| 截断 | 56.8% | **0.9%** | 0.9% | 1.5% |
+| Errors | 18 | 363 (8.2%) | 9 | 0 |
+| Mean | 4.74s | 4.88s | 3.58s | 1.43s |
+| P50 | 0.08s* (假) | 1.75s | 1.52s | 0.65s |
+| P90 | 12.14s | 12.67s | 9.23s | 3.61s |
+| TTFT P50 | - | 0.36s | 0.33s | 0.09s |
+
+✅ **截断从 56.8% 降到 0.9%，路由问题彻底解决**。
+❌ 但 P50 仍然是 baseline 的 2-3 倍。
+
+### Direct-to-D 路径表现优秀（KVC 该有的样子）
+
+按 execution_mode 拆开看：
+
+| 路径 | Exp1 1P7D 占比 | Exp1 1P7D P50 | Exp1 1P7D TTFT P50 |
+|------|:---:|:---:|:---:|
+| `kvcache-direct-to-d-session` ✨ | 42.0% | **0.495s** | **0.043s** |
+| `pd-router-fallback-large-append-session-cap` 🔥 | **52.6%** | 5.6s | 3.7s |
+
+Direct-to-D 路径下：
+- P50 = 0.495s（**比 baseline 0.65s 快 25%**）
+- TTFT P50 = 0.043s（**比 baseline 0.093s 快 2 倍**）
+- KV transfer = 0（无 P 介入，纯 D 上 append-prefill）
+
+这才是 KVC 真正的价值。但只有 30-42% 请求走到这条路。
+
+### 新瓶颈：session-cap fallback 占了 52-65%
+
+`pd-router-fallback-large-append-session-cap` 占 1P7D 的 52.6%、2P6D 的 65.4%。这条路径意味着 router 想开新 session 在 D 上，但 admission 拒绝了（"d-session-cap"），只好回退到 plain router（P 全量 prefill + 传给 D，无 session 复用）。
+
+### Bimodal session 分布（starvation）
+
+| Session | Total turns | Direct-to-D | Session-cap fallback |
+|---------|:---:|:---:|:---:|
+| 22080 | 129 | **98%** | 0% |
+| 3840 | 118 | **97%** | 0% |
+| 70560 | 150 | **0%** | **99%** |
+| 39360 | 148 | **0%** | **99%** |
+| 61600 | 117 | **0%** | **99%** |
+
+要么完全幸运，要么完全饿死——典型的双峰分布。
+
+### 根因：硬编码 cap=4
+
+看 `replay.py:_decode_session_soft_cap` 原始代码：
+
+```python
+def _decode_session_soft_cap(...) -> int:
+    target_tokens = max(1, _estimate_session_resident_tokens(request))
+    usable_capacity_tokens = _usable_capacity_tokens(residency, server_url)
+    ...
+    if usable_capacity_tokens <= 0:
+        return 4
+    return max(1, min(4, usable_capacity_tokens // target_tokens))
+    #              ^^^ 硬编码上限 4
+```
+
+7 个 D × 每个 D 最多 4 个 session = **28 个 session slot 总容量**。Trace 有 52 个 session → 24 个 session 永远抢不到 slot。
+
+启动期 race condition 决定了哪些 session 是"幸运儿"——前 28 个挤进来的 session 的所有后续 turn 都走 direct-to-D（快）；剩下 24 个 session 永远走 session-cap fallback（慢）。
+
+## v4 改进：把硬 cap 从 4 提到 16
+
+`replay.py:_decode_session_soft_cap` 一行修改：
+
+```diff
+-    if usable_capacity_tokens <= 0:
+-        return 4
+-    return max(1, min(4, usable_capacity_tokens // target_tokens))
++    if usable_capacity_tokens <= 0:
++        return 16
++    return max(1, min(16, usable_capacity_tokens // target_tokens))
+```
+
+7 D × 16 = 112 个 slot，远超 52 个 session 需求。
+
+### v4 实际结果（vs v3 1P7D / 2P6D）
+
+| 指标 | v3 1P7D | **v4 1P7D** | v3 2P6D | **v4 2P6D** | baseline 8DP |
+|------|:---:|:---:|:---:|:---:|:---:|
+| Errors | 363 (8%) | 435 (10%) | 9 (0%) | **403 (9%)** | 0 |
+| 截断 | 42 | 43 | 42 | 36 | 68 |
+| **direct-to-D** | 38.6% | **54.3%** | 30.5% | **58.0%** ⭐ | - |
+| **session-cap fallback** | 48.3% | 37.4% | 65.4% | **34.7%** | - |
+| Session reused | 1716 | 2180 | 1358 | **2348** | - |
+| KV transfer blocks | 62K | 53K | 79K | **51K** | - |
+| Mean | 4.88s | 4.21s | 3.58s | **2.51s** | 1.43s |
+| **P50** | 1.75s | 1.08s | 1.52s | **0.84s** | **0.65s** |
+| P90 | 12.67s | 13.38s | 9.23s | **6.51s** | 3.61s |
+| P99 | 28.72s | 24.45s | 18.70s | 18.34s | 8.38s |
+| **TTFT P50** | 0.36s | 0.056s | 0.33s | **0.051s** ⭐ | 0.094s |
+| TTFT P90 | 10.97s | 11.90s | 6.95s | **2.64s** | 0.26s |
+
+✓ direct-to-D 占比从 v3 的 30-38% 涨到 v4 的 54-58%
+✓ session 复用 +27% (1P7D) / +73% (2P6D)
+✓ KV transfer 量 -15% (1P7D) / -36% (2P6D)
+✓ TTFT P50 反超 baseline 46%（0.051s vs 0.094s）
+
+### Direct-to-D 路径全面碾压 baseline（KVC 真实价值）
+
+| Config | n | Lat P50 | Lat P90 | TTFT P50 | TTFT P90 |
+|--------|:---:|:---:|:---:|:---:|:---:|
+| baseline 8DP | 4381 | 0.66s | 3.65s | 0.094s | 0.256s |
+| v4 1P7D direct-to-D | 2179 | 0.495s | 3.03s | 0.044s | 0.055s |
+| **v4 2P6D direct-to-D** | **2348** | **0.499s** | **2.86s** | **0.043s** | **0.054s** |
+
+direct-to-D 子集相对 baseline：
+- P50 快 24-30%
+- P90 快 16-22%
+- TTFT P50 快 54%
+- TTFT P90 快 79%
+
+### 整体性能（去掉 errors 和 truncated）vs baseline
+
+| Config | clean | Mean | P50 | P90 | P99 |
+|--------|:---:|:---:|:---:|:---:|:---:|
+| baseline 8DP | 4381 | 1.45s | 0.66s | 3.65s | 8.38s |
+| v4 2P6D | 4010 | 2.53s | 0.85s | 6.55s | 18.33s |
+
+vs baseline：P50 慢 28%、P90 慢 80%、P99 慢 119%。即使错误率为 0，整体仍输 baseline——根因是 35% 请求被推到 fallback 路径。
+
+### 新瓶颈 1：35% 请求仍走 session-cap fallback
+
+抬到 16 后真实瓶颈是 capacity-based 计算：`min(16, usable_capacity_tokens // target_tokens)`。
+- `target_tokens = input + output`，agentic 里常见 50-100K
+- D 的 KV pool ≈ 100-150K tokens（80GB H100, mem_fraction=0.835）
+- `usable / target` = 1-2，远没到 16 → 真实 cap 是 capacity 算出来的小数字
+
+要解决必须改 capacity-based 估算逻辑（或上方案 D，让 D 自己决定）。
+
+### 新瓶颈 2：9-10% errors（mooncake 传输超时）
+
+P-side log 显示：
+
+```
+KVTransferError: Failed to send kv chunk of <bootstrap_room> to 10.45.7.165:40319
+Sync batch data transfer timeout after 32722558107ns  (32 秒超时)
+Decode instance could be dead, remote mooncake session ... is not alive
+```
+
+特征：
+- 所有 errors 在 run 的 44.8% 之后出现（系统压力累积）
+- 98% errors 集中在 turn ≥ 31（大 input 的请求）
+- v3 cap=4 时 1P7D 已有 363 errors（仅 1 个 D 集中受冲击），v4 cap=16 把压力均匀分布但量级更大
+
+是 mooncake TCP loopback 在并发上去后撞超时，**不是 SGLang 逻辑 bug**。修复方向：
+1. 加长 mooncake transfer timeout（现在 32s）
+2. 限制并发 inflight transfer 数量
+3. 改用 RDMA（loopback 是单机模拟，生产环境换真 RDMA）
+4. chunked KV transfer
+
+## v5 落地方案 D：worker-mode 驱动 seed/reseed
+
+`scripts/sweep_tp1_v5_optD.sh` 真正把方案 D 落到了代码里。改动核心：把 `--kvcache-admission-mode` 从 `local`(replay 估算) 改成 `worker`(D 决策)，并扩展到 **direct_append + seed + reseed 全部路径**。
+
+### 关键代码改动
+
+1. SGLang 侧：`scheduler.py` 的 `admit_direct_append` 端点新增 `mode` 字段，支持 `direct_append | seed`，seed 模式会触发 D 真正去 reserve KV pool 块并主动调用 `maybe_trim_decode_session_cache` 做 LRU。
+2. Replay 侧：`replay.py` 中 reseed / turn-1 seed / large-append-reseed 都改走同一个 admit endpoint；`_decode_session_soft_cap` 在 worker mode 下被完全 bypass。
+3. 新增运行参数：`--kvcache-admission-mode worker`、`--kvcache-seed-min-turn-id 1`、`--kvcache-seed-max-inflight-decode -1`、`--kvcache-prefill-backup-policy release-after-transfer`、`--kvcache-prefill-priority-eviction`。
+
+### 假设
+
+- v4 的 35% session-cap fallback 来自 replay 视图过期 + capacity-based 计算保守 → 让 D 自己看 KV pool 应该把这 35% 救回来。
+- D 主动 LRU eviction 比 replay 自己写的 reservation 更准确，**应该**让更多 session 能 seed 进来。
+
+### v5 实际结果（vs v4 同配置）
+
+| 指标 | v4 1P7D | **v5 1P7D** | v4 2P6D | **v5 2P6D** | baseline 8DP |
+|------|:---:|:---:|:---:|:---:|:---:|
+| Errors | 435 (10%) | **9 (0.2%)** ⭐ | 403 (9%) | **9 (0.2%)** ⭐ | 0 |
+| 截断 | 43 | 42 | 36 | 42 | 68 |
+| direct-to-D | 54.3% | 44.7% ↓ | 58.0% | 41.3% ↓ | - |
+| **session-cap fallback** | 37.4% | **45.6%** ↑ | 34.7% | **50.6%** ↑ | - |
+| no-d-capacity fallback | 0.3% | 1.2% | 0.2% | 0.8% | - |
+| pd-router-turn1-seed (新可见) | - | 1.2% | - | 1.1% | - |
+| pd-router-d-session-reseed (新可见) | - | 4.8% | - | 3.4% | - |
+| pd-router-large-append-reseed (新可见) | - | 1.0% | - | 1.0% | - |
+| Session reused | 2180 | 1990 | 2348 | 1837 | - |
+| KV transfer blocks | 53K | 66K | 51K | 69K | - |
+| Mean | 4.21s | 5.18s | 2.51s | 3.49s | 1.45s |
+| **P50** | 1.08s | 1.59s | 0.84s | 1.31s | 0.66s |
+| P90 | 13.38s | 14.67s | 6.51s | 9.09s | 3.65s |
+| P99 | 24.45s | 26.09s | 18.34s | 24.92s | 8.38s |
+| TTFT P50 | 0.056s | 0.21s | 0.051s | 0.24s | 0.094s |
+| TTFT P90 | 11.90s | 13.06s | 2.64s | 6.90s | 0.26s |
+
+✅ **可靠性大幅提升**：mooncake 传输超时 errors 从 9-10% 跌到 0.2%。D 真容量决策避免了 v4 那种"乐观 admit → 30s 后超时"的死亡链路。
+✅ reseed / turn1-seed 路径首次显式出现，证明 admission 端点对 seed 模式确实生效了。
+❌ **session-cap fallback 不降反升**（37→46% 与 35→51%）。说明 v4 的本地 soft_cap 实际上**比 D 真实容量更乐观**——admit 进来后转身就 OOM，统计成了 error 而不是 fallback。
+❌ 直接结果：**direct-to-D 占比下降、整体延迟全面变差**。P50/P90/P99 与 TTFT 都退步。
+
+### Direct-to-D 子集还是稳的（KVC 真实价值仍在）
+
+| Config | n | Lat P50 | Lat P90 | TTFT P50 | TTFT P90 |
+|--------|:---:|:---:|:---:|:---:|:---:|
+| baseline 8DP | 4381 | 0.66s | 3.65s | 0.094s | 0.256s |
+| v4 2P6D direct-to-D | 2348 | 0.499s | 2.86s | 0.043s | 0.054s |
+| **v5 1P7D direct-to-D** | 1990 | 0.475s | 3.04s | 0.043s | 0.055s |
+| **v5 2P6D direct-to-D** | 1837 | 0.483s | 3.04s | 0.043s | 0.054s |
+
+direct-to-D 的尾延迟和 TTFT 与 v4 几乎完全一致（端点决策开销可忽略），**v5 的回退不是路径本身变慢，而是更多请求被赶到 fallback**。
+
+### Fallback 路径反而比 v4 更糟
+
+| Config | n | Lat P50 | Lat P90 | TTFT P50 |
+|--------|:---:|:---:|:---:|:---:|
+| v5 1P7D session-cap fallback | 2027 | 6.38s | 17.47s | 4.49s |
+| v5 2P6D session-cap fallback | 2253 | 3.13s | 11.25s | 0.89s |
+
+由于 fallback 占比上升、且这条路径本身就比 direct-to-D 慢一个数量级，整体均值被拖累得更厉害。
+
+### v5 真正暴露的瓶颈：D 的 KV pool 物理容量
+
+把 admission 决策权交给 D 之后，瓶颈从"replay 估得太死"变成"D 真的装不下"：
+
+- 80GB H100 × `mem_fraction_static=0.835` → D 单卡 KV pool ≈ 100-150K tokens
+- agentic 长 context session 单 turn footprint 50-100K
+- 单 D 上能并存的 session 数量本就 2-3 个 → 7 个 D 装 50 session 基本不可能
+
+v4 的 cap=16 之所以"看起来好"，部分是因为本地 soft_cap 没真的查 D 的 free pool，开了一堆**最终会失败**的 session（统计成 errors 而非 fallback）。v5 把这部分洗成了"诚实的拒绝"——可靠性跃升的代价是看见了真实容量上限。
+
+### v6 应该针对什么
+
+把 D 物理容量管理打开，而不是再调 replay：
+
+1. **prefill backup 提早 release**（已经加了 `release-after-transfer` 但可能还不够及时） → 让 P 上的 backup blocks 不要长期占用 KV pool。
+2. **priority eviction 策略调优**（已开 `--kvcache-prefill-priority-eviction`）：当前 LRU 可能把 hot session 误踢；需要按 session 命中频率/最近访问做加权。
+3. **chunked / streamed seed**：不要一次 reserve 整个 prompt 的容量，按 chunk 分摊。
+4. **跨 D 的 session migration**：当一个 D 满了但隔壁 D 空时主动迁移，而不是直接 fallback 到 P。
+5. **真正的多机 RDMA**：单机 mooncake loopback 是 errors 的根因之一；上多机 + RDMA 才能让 prefill backup release 后的 KV transfer 真的稳。
+
+工程量：1-3 是 SGLang 内部改 (`scheduler.py` + `session_controller.py`)，4 需要 router 协议扩展，5 是部署变更。
+
+## 关键文件与代码位置索引
+
+| 现象 | 代码位置 |
+|------|----------|
+| Replay policy round-robin | `policies.py:63-67` `RoutingState.next_decode_worker_id` |
+| KV-aware policy（session 亲和） | `policies.py:146-187` `KvAwarePolicy.select` |
+| PD router decode 选择 | `pd_router.py:51-74` `_select_decode_index` |
+| Header 构建 | `replay.py:2407-2424` `_build_headers` |
+| Policy → router config 映射 | `benchmark.py:287-300` `_decode_policy_for/_header_mode_for` |
+| Session admission 软 cap | `replay.py:889-905` `_decode_session_soft_cap` |
+| 已有的 D 侧 admission 端点 | `scheduler.py:3497-3580` `admit_direct_append`（v5 扩展支持 `mode=seed`） |
+| Worker-mode admission 调用方 | `replay.py` reseed / turn1-seed / large-append-reseed 路径 |
+| Prefill backup 释放策略（v5 引入） | `--kvcache-prefill-backup-policy release-after-transfer` |
+| Prefill priority eviction（v5 引入） | `--kvcache-prefill-priority-eviction` |
+| Session 在 D 上找不到的报错 | `scheduler.py:1824-1836` |
+| `_should_allow_local_prefill_on_decode` | `scheduler.py:1975-1980` |
+| Reseed 流程入口 | `replay.py:1665-1809` `_invoke_kvcache_seeded_router` |
+| Direct-to-D 流程 | `replay.py:2351-2398` `_invoke_decode_session_direct` |
+
+## 经验教训
+
+1. **policy 和 mechanism 是两个正交维度**——`--policy default` 不是"无脑默认值"，它真的是 round-robin 无 session 亲和性。KVC 机制必须配 session 亲和的 policy。
+
+2. **不要无脑相信前一个 agent 的 RESULTS_SUMMARY**——v2 的诊断（"local prefill bug"）和实际 finish_reason（"session id does not exist"）完全对不上。任何错误诊断必须用 finish_reason、execution_mode 这些原始字段交叉验证。
+
+3. **bimodal 分布是 starvation 的强信号**——v3 数据里某些 session 100% 走快路径、某些 100% 走慢路径，几乎肯定是某种"先到先得"的资源竞争。看到这种模式立刻去找硬编码 cap 或全局共享资源。
+
+4. **测量要看分组而非整体均值**——v3 整体 P50=1.5s 看似比 baseline 慢，但拆开看 direct-to-D 子集 P50=0.495s 已经反超 baseline。整体均值被 fallback 路径拖累，但 KVC 的核心价值是真实存在的。
+
+5. **errors 与 fallback 是同一类资源压力的两副面孔**——v4 的"低 fallback 率 + 高 error 率"不是更优解，是把容量超限的失败从"显式拒绝"伪装成"超时失败"。v5 把决策权交给真容量后，fallback 升、errors 降，这是更诚实的指标，不要被 v4 的 fallback 数字误导。当看到错误率和 fallback 率呈反相关时，要警惕 admission 决策是否在说谎。

docs/KVC_ROUTER_ALGORITHM.md

@@ -0,0 +1,356 @@
+# KVC-Router：面向 Agentic 多轮 LLM Serving 的 Session-Aware 调度算法
+
+**性质**：论文级形式化规范——用于团队内部对齐 + 外部读者 onboarding。
+**对象**：项目团队（统一术语）；论文 reviewer（算法定义）。
+**最近更新**：2026-05-11。
+
+本文给出本项目所开发的 **KVCache-Centric Router**（以下简称 "KVC-Router"）调度算法的形式化、与实现无关的定义。本文设计为可直接被论文引用，并作为"KVC 到底在谈论什么调度算法"的标准回答。
+
+对应的参考实现位于：
+- `src/agentic_pd_hybrid/policies.py` — `KvAwarePolicy`、`RoutingState`
+- `src/agentic_pd_hybrid/replay.py` — orchestration：admission RPC、reset-on-success、fallback chain
+- `third_party/sglang/python/sglang/srt/managers/scheduler.py` — D-worker 端的 admission 决策
+
+---
+
+## 1. 问题定义
+
+我们要服务一群多轮 agentic LLM session（如 Claude Code、Codex、Cursor 等 coding agent），底层是异构 worker 池，分成：
+- **Prefill workers**（`P`）：GPU 常驻的模型副本，针对长输入 prompt 的 batched prefill 做了优化。
+- **Decode workers**（`D`）：GPU 常驻的模型副本，配备 session-aware KV cache（"SessionAwareCache"），具备：(i) 跨 turn 保留 session 的 KV 状态；(ii) 在本地已缓存的 prefix 上做 append-prefill，无需绕回 `P`。
+
+在一个 agent turn 内，请求 `r` 到达时其对话 prefix 已经从前序 turn 累积；**新增**的 tokens（工具输出、用户消息等）构成小规模 **append**。驱动 KVC 设计的根本观察是：
+
+> 当 prefix KV **已经驻留在将要解码该请求的 D worker 上**，请求的 first-token 延迟仅由 *append* 大小决定（典型 O(10²–10³) tokens），而非完整 prompt 大小（典型 O(10⁴–10⁵) tokens）。
+
+Router 的工作就是最大化满足上述条件的请求占比，同时尊重容量约束、不造成 session 无限饿死。
+
+### 1.1 优化目标
+
+给定来自 `S` 个 session 的请求流 `R = (r_1, r_2, ...)`，最小化 SLO 加权的 TTFT 与端到端延迟混合：
+
+```
+   minimize   E[ w_ttft · TTFT(r) + w_lat · E2E_Latency(r) ]
+   subject to  capacity[d] ≤ K_d   对任意 D worker d 在任意时刻 t,
+               没有 session 被永久拒绝服务.
+```
+
+参考实现中通过 measurement 隐式取 `w_ttft = 1, w_lat = 1`；per-D KV 池预算 `K_d` 取 SGLang 启动时上报的 `max_total_num_tokens`。
+
+---
+
+## 2. 系统模型与记号
+
+### 2.1 集合
+
+| 符号 | 含义 |
+|---|---|
+| `P = {p₁, …, p_|P|}` | Prefill worker 池 |
+| `D = {d₁, …, d_|D|}` | Decode worker 池 |
+| `S` | Session 标识符集合（由上游 agent runtime 分配） |
+| `H` | KV block hash 的全集（本实现中每 `BLOCK_TOKEN_BUDGET = 24` tokens 对应一个 hash） |
+
+### 2.2 请求
+
+一个请求 `r` 是一个元组：
+
+```
+   r = ⟨ s(r),  t(r),  prefix_hashes(r),  append_len(r),  input_len(r) ⟩
+```
+
+其中：
+- `s(r) ∈ S` — session id
+- `t(r) ∈ ℕ` — 该 session 内的 turn index（0 = 首轮）
+- `prefix_hashes(r) ⊂ H` — 覆盖请求输入 prefix 的 block hash 集合
+- `append_len(r) ∈ ℕ` — 新到达、**不在** `prefix_hashes(r)` 中的 token 数
+- `input_len(r) = (|prefix_hashes(r)| · 24) + append_len(r)` — 总 token 数
+
+### 2.3 Router 状态 (`Σ`)
+
+Router 跨请求维护的全局状态：
+
+| 字段 | 类型 | 语义 |
+|---|---|---|
+| `resident[d]` | `set[H]` | Router 估计的 D `d` 当前 SessionAwareCache 中常驻的 block hash 集合（router 端估计，真值在 worker 上） |
+| `pin[s]` | `D ∪ {⊥}` | Session `s` 最近一次成功服务的 D；`⊥` 表示从未见过 |
+| `inflight[d]` | `ℕ` | 当前已派发给 `d` 但尚未完成的请求数 |
+| `assigned[d]` | `ℕ` | 累计派发到 `d` 的路由决策次数（负载 tie-breaker） |
+| `rejects[s,d]` | `ℕ` | per-(session, D) 的 admission 拒绝计数（v2 引入的 migration 机制） |
+
+### 2.4 超参数
+
+| 符号 | 默认值 | 描述 |
+|---|---|---|
+| `α`（`sticky_bonus`） | 1 | 匹配 `pin[s]` 的 D 在评分中获得的 bonus |
+| `τ_reject`（`migration_reject_threshold`） | 3 | (s, d) 被拒绝达此次数后，d 对 s 进入 blacklist |
+| `τ_append`（`kvcache_direct_max_uncached_tokens`） | 8192（v2） | 走 Direct-to-D 路径允许的最大 append 长度 |
+| `K_d` | 取自 SGLang `max_total_num_tokens` | per-D 的 KV 池预算 |
+| `ρ` | 0.95 | 容量高水位线（隐式由 SGLang 强制） |
+| `ε`（最大 fallback 重试数） | `|D| - 1` | router 在退化到 vanilla PD-disagg 之前最多探测几个 D |
+
+### 2.5 路由结果
+
+路由决策 `δ(r)` 取以下四种之一：
+
+| Mode | 含义 | KV transfer |
+|---|---|---|
+| `Direct(d)` | r 完全在 D `d` 上执行；D 在其常驻 KV 上做 append | **无**（快路径） |
+| `Seed(d)` | Session 首轮：P 做完整 prefill，KV 通过 mooncake 传到 `d` | 完整 input |
+| `Reseed(d)` | Session 之前在某个 D' 上，但已不再常驻；按 Seed 处理 | 完整 input |
+| `Fallback(p, d)` | Vanilla pd-disagg 路径（其它 D 均被 blacklist 或拒绝） | 完整 input |
+
+---
+
+## 3. 算法
+
+KVC-Router 由三个相互配合的过程组成：
+- **Algorithm 1 (`Route`)**：router 端基于评分的候选选择。
+- **Algorithm 2 (`Admit`)**：D-worker 端的 admission 决策（在 D scheduler 中执行，非 router）。
+- **Algorithm 3 (`Dispatch`)**：端到端 orchestration，把 Route + Admit + reset-on-success 串起来。
+
+### 3.1 Algorithm 1：`Route(r, Σ)` — 基于评分的候选选择
+
+```
+输入：请求 r，状态 Σ
+输出：候选 d* ∈ D（若所有 D 都被过滤后仍无候选，退化分支兜底返回最少被拒的 D）
+
+ 1.  blacklisted ← { d ∈ D : Σ.rejects[s(r), d] ≥ τ_reject }
+ 2.  C ← D ∖ blacklisted                                  // 候选 D 集合
+ 3.  if C = ∅ :                                           // 退化
+ 4.       return argmin_{d ∈ D} Σ.rejects[s(r), d]        // 选最少被拒的 D
+ 5.  for each d ∈ C :
+ 6.       overlap(d)  ← |prefix_hashes(r) ∩ Σ.resident[d]|
+ 7.       sticky(d)   ← 1 if Σ.pin[s(r)] = d else 0
+ 8.       infl(d)     ← Σ.inflight[d]
+ 9.       assn(d)     ← Σ.assigned[d]
+10.       score(d)    ← ⟨ overlap(d) + α·sticky(d),       // 主项
+                          sticky(d),                       // tie-1
+                          −infl(d),                        // tie-2（负载小者占优）
+                          −assn(d) ⟩                       // tie-3
+11.  return argmax_{d ∈ C} score(d)                       // 按字典序最大
+```
+
+**说明**：
+- 评分是 **4 元组按字典序比较**，不是单个标量——这样避免在不同维度之间调权重。
+- 第 10 行的主项 `overlap + α·sticky` 同时奖励 KV 复用与 session stickiness。取 `α=1`、`overlap` 以 block（24 tokens）为单位时，**任何一次 hash 命中都压制纯 sticky 的候选**。
+- 第 1–4 行的 blacklist 过滤防止永久绑死在已饱和的 D 上；与 Algorithm 3 的 reset-on-success 配合，限定了 migration 频率。
+
+### 3.2 Algorithm 2：`Admit(d, r, M, K)` — D-worker admission 决策
+
+在 D worker 自己的 scheduler 内部执行（非 router），这是 **KVC 的机制核心**：每个 D 自治判断能否把 `r` 当作 Direct（append-only）服务，还是必须改走 P 路径。
+
+```
+输入：D worker d，请求 r，d 上本地常驻的 session 集合 M_d，KV 池预算 K_d
+输出：⟨can_admit ∈ {True, False},  mode ∈ {Direct, Seed, Reseed, ⊥},  reason⟩
+
+ 1.  used_tokens ← Σ_{s' ∈ M_d} resident_tokens(s', d)     // D 自己的 bookkeeping
+ 2.  cap_ok ← (used_tokens + input_len(r)) ≤ ρ · K_d        // 高水位线 ρ ≈ 0.95
+
+ 3.  if s(r) ∈ M_d :                                        // session 在 d 上有常驻
+ 4.       if append_len(r) ≤ τ_append  and  cap_ok :
+ 5.           return ⟨True, Direct, ∅⟩                      // → 快路径
+ 6.       elif append_len(r) > τ_append :
+ 7.           return ⟨False, ⊥, "real-large-append"⟩
+ 8.       else :
+ 9.           return ⟨False, ⊥, "no-d-capacity"⟩
+
+10.  else :                                                 // session 在 d 上无常驻
+11.       if cap_ok :
+12.           mode ← Seed if t(r) = 0 else Reseed
+13.           return ⟨True, mode, ∅⟩                        // → 经 P 做 KV seeding
+14.       else :
+15.           return ⟨False, ⊥, "session-not-resident-no-capacity"⟩
+```
+
+**说明**：
+- 该过程通过同步 HTTP RPC（`/admit_direct_append`）从 router 调用。RPC 阻塞直到 D scheduler 给出权威答复——这是 v5 引入的 **"worker-mode admission"**，替换了更早的 router-端容量估算（系统性偏乐观）。
+- reason 字符串被回传给 router，用于：(i) 在 Algorithm 3 中驱动 fallback chain；(ii) 标注 `execution_mode` 字段便于分析。
+
+### 3.3 Algorithm 3：`Dispatch(r, Σ)` — 端到端 orchestration
+
+```
+输入：请求 r，状态 Σ
+输出：执行模式 μ ∈ {Direct, Seed, Reseed, Fallback}
+
+ 1.  retries ← 0
+ 2.  tried ← ∅
+ 3.  while retries < ε :
+ 4.       d* ← Route(r, Σ \ {对 tried 中的 d 已 bump 过的 rejects})
+ 5.       if d* = ⊥ : break                                  // 无候选
+ 6.       resp ← Admit(d*, r)                                // RPC 到 D scheduler
+ 7.       if resp.can_admit :
+ 8.           Σ.rejects[s(r), d*]  ← 0                       // ◀ reset-on-success（v2）
+ 9.           Σ.pin[s(r)]          ← d*
+10.           Σ.inflight[d*]       ← Σ.inflight[d*] + 1
+11.           if resp.mode = Direct :
+12.                在 d* 上完整执行 r（append-prefill + decode）
+13.                return Direct
+14.           else :                                          // Seed 或 Reseed
+15.                p ← round_robin_next(Σ, P)
+16.                在 p 上做 r 的 prefill
+17.                经 mooncake 把 KV(r) 从 p 传到 d*
+18.                在 d* 上 decode r
+19.                return resp.mode
+20.       else :
+21.           Σ.rejects[s(r), d*]  ← Σ.rejects[s(r), d*] + 1
+22.           tried ← tried ∪ {d*}
+23.           retries ← retries + 1
+24.
+25.  // ε 次重试耗尽——退化 Fallback 到 vanilla pd-disagg
+26.  p ← round_robin_next(Σ, P)
+27.  d ← round_robin_next(Σ, D)
+28.  通过 ⟨p, d⟩ 走 pd-disagg(r)
+29.  return Fallback
+```
+
+**维持的关键不变量**：
+
+1. **不会静默过载**：一个 D 永不接受会让 `used_tokens > ρ · K_d` 的请求（Algorithm 2 第 2 行）。
+2. **不存在永久饿死**：对任意 session `s`，只要曾在某 D `d*` 上成功过一次，之后 `Σ.rejects[s, d*] = 0`（Algorithm 3 第 8 行）。因此 blacklist 计数器不会对仍在某处成功获得服务的 session 累积——这阻止了 **v1 的 thrashing 病理**：原本 blacklist 计数器单调增长 + 退化 fallback 形成自放大的 round-robin 死循环。
+3. **migration 有界**：一个 session 从 D `a` 迁移到 D `b` 必须经过连续 `τ_reject` 次在 `a` 上失败、期间无任何成功。每个 session 生命周期内的最坏 migration 次数 ≤ `(|D| − 1) · τ_reject`。
+
+### 3.4 Reset-on-success：为什么这是关键修复（v1 → v2 演化）
+
+v1 实现**省略了** Algorithm 3 第 8 行——一旦 `(s, d)` 累积 `τ_reject` 次拒绝，d 对该 session **整个 run 永久 blacklist**。实测（Migration v1，见 `docs/MIGRATION_V1_FINDINGS_ZH.md`）触发了自放大的失效模式：
+
+```
+session s 在 d 上稳定服务 70 个 turn
+       ↓ 瞬时 burst 让 d 短暂饱和
+3 次到 d 的 admission 被拒 → rejects[s,d] = 3 → d 对 s 永久 blacklist
+       ↓ s 迁到 d'，d' 也在负载中 → 被拒 → blacklist
+       ↓ d'' 同理
+所有 D 都 blacklist → 退化 fallback round-robin → 每次重试都 bump 一次计数器
+                   → s 永远在 D 之间 thrashing，每次都丢失 KV residency
+```
+
+reset-on-success 关上了这个回路：只要 `s` 在任一 d 上真正完成一次 Direct，针对该 session 的 blacklist 立刻清零。该机制只对**持续性**（不是瞬时性）容量压力触发。
+
+---
+
+## 4. 性质
+
+### 4.1 Theorem 1（在有界 ε 下无永久饿死）
+
+*假设 `τ_reject ≥ 1` 且每个 D worker 的容量非零。则对任意能在 admission 时容下的 session `s`，Algorithm 3 在至多 `|D| · τ_reject` 次重试内返回 `{Direct, Seed, Reseed}` 之一；之后任意一次 Direct 成功即可清空 `s` 的所有 blacklist。*
+
+**证明概要**：每次循环要么成功（return）、要么恰好让某个 `rejects[s, d]` 计数器 +1（第 21 行）。经过 `|D| · τ_reject` 次迭代后，每个 D 要么对 `s` 已被 blacklist（`Route` 第 1 行会过滤），要么已成功（已终止）。在所有 D 都被 blacklist 的饱和点，`Route` 第 3 行返回最少被拒的 D，打破对称性，强制取得进展。∎
+
+### 4.2 Theorem 2（fast-path 命中下限）
+
+*假设 session `s` 在 D `d` 上已积累 KV residency `R_s ⊂ H`，且在某 turn `t > 0` 提交的请求 `r` 满足 `prefix_hashes(r) ⊆ R_s`、`append_len(r) ≤ τ_append` 且 admission 容量充足。则 Algorithm 3 将 `r` 路由为 Direct(d)。*
+
+**证明概要**：由 Algorithm 1，`overlap(d) = |R_s|` 取得最大值；结合 `α·sticky(d) ≥ 1`，d 的字典序得分严格高于任何 `prefix_hashes(r) ⊈ R_{s,d'}` 的 d'。故 `Route` 返回 d。`Admit(d, r)` 进入 `s ∈ M_d ∧ append ≤ τ_append ∧ cap_ok` 分支，返回 Direct。∎
+
+这是 **支持架构设计的机制级保证**：只要 residency、append 大小、容量三者同时成立，快路径就被**确定性地**选中；KVC 在典型场景下的 TTFT 优势是结构性属性，不是概率性。
+
+### 4.3 复杂度
+
+每个请求：
+- `Route`：`O(|D|)`（每个候选 D 算一次 score）。生产规模下 `|D| ≤ 8`，主要开销在 Python 层，≪ 1 ms。
+- `Admit`：D scheduler 内部 O(1)（查自己的 bookkeeping，无全局锁）。
+- Router 层的单请求总开销：`O(|D|)` 计算 + 1 次到目标 D 的 HTTP RTT（loopback 亚毫秒，跨机数据中心约 1 ms）。
+
+---
+
+## 5. 与 baseline 的对比
+
+| 性质 | Vanilla pd-disagg | DP（cache-aware） | **KVC-Router**（本文） |
+|---|---|---|---|
+| P/D 分离 | 是（`|P| + |D|` GPU） | 否（每个 worker fused P+D） | 是 |
+| 跨 turn cache locality | 无（每个请求都 P→D 传 KV） | 仅在单 fused worker 内部走 hash prefix 路由 | session 钉在某 D 上，本地 append-prefill |
+| 同 session cache 集中度 | 无 | 散到 `|D|` 个 worker（每个占 1/|D|） | 集中在一个 D（整段常驻） |
+| 最坏 turn-2 prefill 工作量 | 完整 input 经 P→mooncake→D | 在目标 worker 上做完整 prefill（带 prefix cache 命中） | 本地 `append_len ≤ τ_append` tokens |
+| 容量感知 admission | 无（router 盲发） | 隐式靠 worker 队列深度 | 显式的 per-D `Admit()` 决策 |
+| Migration 机制 | N/A | N/A | 带 reset-on-success 的 reject-counter blacklist |
+| Idle prefill 成本 | 是——P 永远在算 | 否 | 是——P 只在 cache miss 时启用（本工作 SWE-Bench 评测下约 8% 请求） |
+
+KVC 的关键架构权衡：**用 P 端 GPU 闲置换 D 端 TTFT 稳定性**。在 per-session cache 复用率高的 agentic workload 上（Inferact 的 Codex trace 报告 94.2% cache hit；我们的 SWE-Bench replay 实测 91.6% Direct 命中），这个交换显著有利。在 session 短或 cache hit 低的 workload 上，权衡反转、DP 胜出。
+
+---
+
+## 6. 符号速查表
+
+| 符号 | 含义 |
+|---|---|
+| `P, D` | Prefill / Decode worker 池 |
+| `s(r), t(r)` | 请求 r 的 session id 与 turn index |
+| `prefix_hashes(r)` | r 输入 prefix 的 KV block hash |
+| `append_len(r)` | r 中新增（未缓存）部分的 token 数 |
+| `Σ.resident[d]` | Router 对 d 缓存 block 集合的估计 |
+| `Σ.pin[s]` | session s 最近一次成功的 D |
+| `Σ.rejects[s,d]` | per-(s,d) 的 admission 拒绝计数 |
+| `α` | sticky bonus 权重（默认 1） |
+| `τ_reject` | migration 阈值（默认 3） |
+| `τ_append` | Direct 路径允许的 max append 大小（v2 默认 8192） |
+| `K_d` | D worker d 的 KV 池预算 |
+| `ρ` | 容量高水位（默认 0.95） |
+| `ε` | fallback 重试上限（默认 `|D| − 1`） |
+| `δ(r)` | 路由决策：`Direct(d)` / `Seed(d)` / `Reseed(d)` / `Fallback(p, d)` |
+
+---
+
+## 7. 本工作评测中实际使用的默认参数
+
+| 参数 | 取值 | 说明 |
+|---|---|---|
+| `|P|, |D|` | 1, 3（1P3D 配置） | 单机 4× H100 80GB |
+| `α` | 1 | |
+| `τ_reject` | 3 | |
+| `τ_append` | 8192 | v2 调优后取值（v0/v1 用 2048） |
+| `K_d` | 92104 tokens | SGLang 按 `mem_fraction_static=0.835` 自动算出 |
+| `ρ` | 隐式 ~0.95 | 由 SGLang 的 `max_total_num_tokens` 强制 |
+| `ε` | 2 | `|D| − 1 = 2` |
+| 每次 run 的 session 数 | 52 | SWE-Bench 50sess trace |
+| 总请求数 | 4449 | |
+| Time-scale | 1.0（真实 trace 时序） | |
+| 并发 | 32 | |
+
+---
+
+## 8. Anti-patterns（KVC **不**是什么）
+
+1. **KVC 不仅仅是 kv-aware routing**。DP 和 KVC 都可以跑 `kv-aware` policy；KVC 在此之上加了三件事：(i) session 钉定，(ii) worker 端 admission，(iii) 带 reset-on-success 的 migration。如果在比较 "KVC vs DP" 时缺这三个要素的任何一个，**测的就不是 KVC 与 DP 的差异**。
+
+2. **KVC 在 policy 项里不直接感知容量**。`Route` 不查 per-D 容量；容量感知完全经由 `Admit` 拒绝来传导。我们刻意做了这层分层——把容量判断放进 `Route` 会引入"换 D"的决策空间，导致 orphan KV 滞留问题。
+
+3. **KVC 不保证 load balance**。一个 session 若能舒服地装在某个 D 上，可能永远钉在那里，而其它 D 大部分时间空闲。在低容量压力下这是设计意图；高压力下 Theorem 1 的 migration 会触发再均衡。
+
+4. **`Fallback` 不是"降级路径"**。它和 vanilla pd-disagg 请求结构性等价，延迟特征相同。KVC 的价值在于让 Fallback 占比在典型 agentic workload 下 ≪ 10%。
+
+---
+
+## 9. 公开问题（reviewer 关注点）
+
+以下问题在当前评测中尚未解决，主动列出以保持透明：
+
+1. **Session 钉定相对于纯 P/D disaggregation 的边际贡献是多少？** 需要 `naive 1P3D` 对照实验（vanilla SGLang xPyD，不带 KVC 层）——仓库当前缺失（见 `docs/V2_DEEP_ANALYSIS_ZH.md §4.7`）。
+
+2. **Algorithm 3 在更高压下行为如何**（例如 ts=10 加速、session 数 ≫ |D|·K_d/peak_input）？当前 ts=1 评测对应真实 agentic 区间，但算法在更高负载下的鲁棒性未经实验验证。
+
+3. **真 RDMA 下的 reseed 代价**：本次评测的 3–7 s reseed 延迟由两段组成——P 端 re-prefill（1.5-3s）+ P→D mooncake transfer（1.5-4s）。当前 sweep 用的是 TCP loopback；启用 IB/RoCE（节点有 mlx5_0/_1 @ 200 Gb/s × 2 active，需在 sweep 加 `--force-rdma --ib-device mlx5_0`）只能压缩 transfer 段到 ~200ms，**不动 re-prefill 段**。预期 TTFT p99 从 1.28s 降到 ~0.7s（仍输 DP 0.43s）。待独立验证。
+
+4. **D→P 增量 KV 同步（核心 future-work 缺口）**：reseed 长尾的真正消除需要让 P 端 backup 跟上 D 的 direct-to-D append 增长。经独立 forensic 审查，**当前代码、vendored SGLang、mooncake 三层均无 D→P KV transfer 实现**：mooncake `MooncakeKVManager` 是 PREFILL=sender / DECODE=receiver 的硬角色分支（`add_transfer_request` 上有 `assert disaggregation_mode == PREFILL` 硬约束），`BaseKVSender` / `BaseKVReceiver` 抽象无 bidirectional slot，`session_aware_cache.release_session` 在驱逐时只调 `kv_pool_allocator.free()` 无出站，`_commit_prefill_backup_residency` 唯一 caller 是 seed/reseed 路径；`capacity-backup` policy 的真实语义只是"reseed 完不关 P streaming session"——backup 是 seed-time 的静态快照，不随 direct-to-D append 同步。要实现 D→P 增量同步，工程量 ~1-2 周，最难的不是 mooncake 加 D-sender / P-receiver 角色（~400 LOC），而是 **SGLang radix tree 改成允许从外部 worker 喂数据**——radix cache 当前假设单一生产者（本 worker model 输出）。这是论文里最值得做的 contribution 之一。
+
+5. **v2 代码路径下的确定性**：v0 代码库的 ts=1 N=3 categorical 确定性已经证实；新增的 reset-on-success 分支与 threshold=8192 路径未被独立 re-validate。两个额外的 N=1 run 即可解决。
+
+---
+
+## 10. 论文引用建议
+
+论文中提到本算法时建议表述：
+
+> "We use the KVC-Router scheduling algorithm (Algorithms 1–3 of [our paper], formally defined in our supplementary materials). The router selects a decode worker by lexicographic scoring on `(overlap+α·sticky, sticky, −inflight, −assigned)` (Algorithm 1), defers the admission decision to the chosen worker via a synchronous RPC (Algorithm 2), and maintains a per-(session, decode worker) rejection counter that is reset on every successful Direct admission (Algorithm 3). This last detail — reset-on-success — is what distinguishes our v2 from the unstable v1 implementation that exhibits self-amplifying session thrashing."
+
+---
+
+**附录 A — 算法步骤到代码实现的对照**
+
+| 算法步骤 | 文件 | 符号 |
+|---|---|---|
+| `Route` 第 5–11 行 | `policies.py:189–202` | `KvAwarePolicy.select` 内层循环 |
+| `Route` 第 1–4 行（blacklist 过滤 + 退化分支） | `policies.py:182–187, 204–211` | `migration_reject_threshold`，`select` 的 fallback |
+| `Admit` | `third_party/sglang/python/sglang/srt/managers/scheduler.py` | `handle_admit_direct_append_request` |
+| `Dispatch` 第 8 行（reset-on-success） | `replay.py: _run_request` | finish 路径中的 reset |
+| `Dispatch` 第 21 行（记录 reject） | `replay.py: _run_request` | `state.record_admission_reject(...)` |
+| 超参数 `τ_append` | CLI flag | `--kvcache-direct-max-uncached-tokens` |
+| 超参数 `τ_reject` | CLI flag | `--kvcache-migration-reject-threshold` |

docs/MIGRATION_V1_FINDINGS_ZH.md

@@ -0,0 +1,283 @@
+# Migration v1 实验发现：blacklist 永久性导致 thrashing
+
+**日期**：2026-05-08
+**状态**：v1 run 进行中（~23% 完成时的中期分析）
+**前置文档**：
+- `docs/REFACTOR_PLAN_V1_ZH.md` §6.2（v1 设计）
+- `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md` §2.1（§1 starvation claim）
+
+**触发**：v1 实现的 session migration（rejection blacklist 机制）部署后，观测到 session-level thrashing——某些 session 在 3 个 D 之间 round-robin 高达 75-116 次。本文记录中期数据、根因诊断、v2 设计。
+
+---
+
+## 0. TL;DR
+
+1. **v1 修复了 §1 starvation 但引入了新的 thrashing 失效模式**——不是 admission 过严，是 blacklist 永久累积的设计 bug
+2. **核心证据**：session 6880 在 decode-1 上稳定 70 turns，然后某瞬时 burst 把 reject 计数累积到阈值，被永久 blacklist，之后陷入 3-D 间 round-robin 死循环
+3. **85% admission 拒绝是 `session-not-resident`**——非 D 真容量问题，而是迁移后"新 D 第一次见你"的正常语义
+4. **v2 设计**：reset-on-success 让 reject 计数在成功 turn 后清零，只有**持续**失败才迁移
+5. **深层观察**：baseline 的"100% pin 但稳定"可能比"分布均匀但 thrashing"更好——糟糕的优化可能比不优化还糟
+
+---
+
+## 1. v1 实施回顾
+
+### 1.1 改动文件
+- `src/agentic_pd_hybrid/policies.py`：`RoutingState.session_d_rejects` Counter；`KvAwarePolicy.migration_reject_threshold=3` skip blacklisted D；degenerate fallback 选最少拒的 D
+- `src/agentic_pd_hybrid/replay.py`：`_run_request` 末尾 `state.record_admission_reject(sess, D)`（基于 execution_mode 子串匹配）；`_fallthrough_reason` 把 `pd-router-fallback-large-append-*` 拆成 `session-not-resident` / `real-large-append` / 等
+- CLI / benchmark wiring
+
+### 1.2 v1 假设（事后看部分错误）
+- "reject 计数 + 阈值 3 = 容忍短期波动 + 持续失败迁移" ← **错**，counter 永久增长导致迁移成必然
+- "迁移到新 D 后 session 在新 D 稳定下来" ← **部分错**，迁移到的新 D 也很可能很快 reject
+- "session-not-resident 不会触发计数" ← **大致对**，但下游 fallback 可能间接触发
+
+---
+
+## 2. 中期数据（1023/4449 reqs，~23%）
+
+### 2.1 头部指标 vs baseline
+
+| 指标 | baseline kvc_1p3d_run1 | v1（中期） |
+|---|---:|---:|
+| Per-D 调用分布 | 1502/1445/1502（±3.8%）| 796/785/779（**±1.1%**，更均衡）|
+| Per-D 峰值 token_usage | 0.99/0.99/0.99 | 0.31/0.30/0.00（**容量充裕**，未顶到 1.00）|
+| KVTransferError | 5（全程）| 6（中期，趋势相近）|
+| 已见 sessions | 52（全程）| 29（中期）|
+
+**好的方面**：
+- 负载均衡度跃升（±26%→±1.1% if normalized）
+- D 容量从未饱和——§2 假设的"D drain time"机制配合 ts=1 充分发挥
+- 0 sessions 永久 stuck 在饿死状态
+
+### 2.2 Migration 触发情况（已见 29 sessions）
+
+| 类别 | 数量 | 占比 |
+|---|---:|---:|
+| 仍 pin 在 1 个 D | 9 | 31% |
+| 触碰 2 个 D | 3 | 10% |
+| **触碰所有 3 个 D** | **17** | **59%** |
+
+**D-切换次数分布**：
+- mean = 26 次/session
+- median = 16 次
+- **max = 116 次**
+- 15 sessions 切换 >10 次（明显 thrashing）
+- **6 sessions 切换 >50 次**（严重 thrashing）
+
+---
+
+## 3. 根因诊断：session 6880 的轨迹
+
+### 3.1 数据
+
+```
+turn 0-70:   全部在 decode-1   (71-turn 稳定 streak)  ← §1 baseline 行为
+turn 71-150: 在 3 个 D 间剧烈 thrashing
+              decode-0: 26 个短 streak
+              decode-1: 25 个短 streak
+              decode-2: 25 个短 streak
+              平均 streak 长度 = 2 turns
+              total streaks = 76
+```
+
+### 3.2 解读
+
+**前 70 turn 完美稳定**：session 6880 在 decode-1 上正常运行 70 个 turn，每次都成功，是 baseline §1 "100% pin" 的复现——稳定但不公平（其他 session 没分到 decode-1 的资源）。
+
+**第 71 turn 后崩溃**：
+1. 某个瞬时 burst（其他 session 的活动？）让 decode-1 短暂饱和
+2. session 6880 在 decode-1 上连续 3 次被 admission 拒（`no-space` 或 `d-session-cap`）
+3. v1 的 `state.session_d_rejects[(6880, decode-1)]` 累积到 3 → blacklist
+4. policy 改选 decode-0 → 同样发生 → blacklist
+5. 改选 decode-2 → 同样 → blacklist
+6. **3 D 全部 blacklisted** → degenerate fallback 在 3 D 间 round-robin
+7. 每次 round-robin 又触发新 reject → 计数继续涨 → 永远在 thrashing 死循环
+
+### 3.3 admission 数据交叉验证
+
+中期 1932 admission events 解构：
+
+| mode × can_admit × reason | count |
+|---|---:|
+| `direct_append, True, None` | 1721（成功）|
+| `direct_append, False, session-not-resident` | **62** |
+| `seed, True, None` | 142（成功）|
+| `seed, False, no-space` | **11** |
+
+**只有 11 个 "no-space" 才是真容量拒绝**（占总 admission 的 0.6%）。62 个 "session-not-resident" 是迁移后"新 D 第一次见你"的正常语义。
+
+但因为 v1 用 `_is_admission_rejection_mode` 通过 execution_mode 子串匹配，下游 fallback chain 会把 `session-not-resident` 也间接累积到计数器（fallback 链路本身可能触发 session-cap）。
+
+---
+
+## 4. 设计 bug 三层
+
+### 4.1 Bug 1：blacklist 永久性
+
+```python
+# policies.py 当前实现
+if rejects >= self.migration_reject_threshold:
+    continue  # skip this D forever
+```
+
+`session_d_rejects[(sess, D)]` 是单调递增 Counter。一旦达到阈值，**永远**被 skip。但 D 的容量是动态的——70 个 turn 后短暂饱和不代表它后续不能服务这个 session。
+
+### 4.2 Bug 2：degenerate fallback 加剧问题
+
+当所有 D 都被 blacklist：
+```python
+best_decode_worker_id = min(
+    (w.worker_id for w in topology.route_workers),
+    key=lambda wid: state.session_d_rejects.get((sess, wid), 0),
+)
+```
+选"最少被拒"的 D。但每次 fallback 又增加该 D 的计数 → 下次选另一个 D → 形成完美 round-robin，永远走不出 thrashing。
+
+### 4.3 Bug 3：信号归并粗糙
+
+`_is_admission_rejection_mode` 子串匹配 `session-cap` / `no-d-capacity` / `d-backpressure`，但执行链路可能这样：
+
+```
+direct_append → session-not-resident（85% 占比，正常迁移后语义）
+  → fallback 试 seed
+    → seed admit ok（142/153 = 93%）→ execution_mode = pd-router-d-session-reseed-*（不计 reject）
+    → seed no-space（11/153 = 7%）→ execution_mode = pd-router-fallback-X-no-d-capacity（计 reject）
+```
+
+绝大多数 fallback 不会触发 reject 计数。但 thrashing 一旦开始，很容易踩到那 7% no-space 路径，calculator 增长一次。15+ 次 thrashing 后，单 D 计数累到 3 完全可能。
+
+**所以设计 bug 不在信号粗糙，而在永久累积 + degenerate round-robin。**
+
+---
+
+## 5. 深层观察：稳定 vs 公平的 trade-off
+
+| | baseline（v0）| v1 |
+|---|---|---|
+| 公平性 | 18/52 永久饿死 | 0 永久饿死 |
+| 稳定性 | 100% pin（结构稳定）| 6/29 严重 thrashing |
+| Per-D 负载均衡 | ±26% | ±1.1% |
+| 大 session 体验 | 慢但稳定（每 turn 都走 fallback ~1.0s）| 不稳定 + 频繁 D 切换 + 丢 KV state |
+
+**预想反直觉的结果**：v1 在头部指标（per-D 均衡）赢，但在 session 体验可能输——
+- baseline 的 fallback 路径有稳定 ~1s latency
+- v1 的 thrashing session 每次 D 切换都 close 旧 session、丢 KV、新 D 上重新建立——有可能 latency 反而更高
+
+需要等 run 结束的 lat mean / TTFT mean 数据验证。**糟糕的优化可能比不优化还糟。**
+
+---
+
+## 6. v2 设计
+
+按 ROI 排序的修复层。**先做 #1，验证后再决定是否需要 #2/#3**。
+
+### 6.1 v2-fix-1：reset-on-success（最高 ROI）
+
+```python
+# replay.py _run_request 末尾，在 state.finish 后
+if execution.execution_mode == "kvcache-direct-to-d-session":
+    # 这次 direct-to-D 成功 = D-X 仍能服务这个 session
+    # 清零累积的 reject 计数（消除永久 blacklist）
+    state.session_d_rejects[(request.session_id, decision.decode_worker_id)] = 0
+```
+
+**预测效果**：
+- session 6880 在 decode-1 上 70 个成功 turn 把计数反复清零
+- 即使中间出现 1-2 次瞬时 reject，下次成功立刻清零
+- 只有**持续**失败（reject 后 reject 后 reject，没有夹杂 success）才能累到阈值
+- 真饿死的 session（如 35680/39360 input >92K）才会触发迁移
+
+**工程量**：~5 行代码 + 1 个 smoke + 1 个完整 run（~5.5h）
+
+### 6.2 v2-fix-2：sliding window（如果 #1 不够）
+
+把 `Counter` 改成 `dict[(sess, D), deque[float]]` 存最近 K 次拒绝时间戳。判断时用最近 N 秒（或 N 个 turn）内的次数。
+
+更稳健但更复杂。**若 #1 已能彻底解决 thrashing，跳过此项。**
+
+### 6.3 v2-fix-3：reject 类型分离（如果 #1 + #2 不够）
+
+把 admission reason 显式传到 _run_request，区分：
+- `no-space` / `session-cap` / `backpressure` → 计 reject
+- `session-not-resident` → 不计
+
+需改 `ExecutionResult` 加 `admission_reject_reason` 字段，并在 fallback 链路传递。**不在第一轮**——先看 #1 是否够用。
+
+### 6.4 v2 应保留的 v1 设计
+
+- 阈值 3（不变）
+- `record_admission_reject` 的子串匹配（不变）
+- 新 fallback labels（`session-not-resident` 等）（不变）
+- degenerate fallback 选最少拒的 D（不变，但因为 reset-on-success 几乎不会触发到此分支）
+
+---
+
+## 7. 实验计划
+
+| 阶段 | 动作 | 时间 |
+|---|---|---|
+| 1 | 等 v1 run 完成（ETA ~16:30）| 自然 |
+| 2 | 跑 analyzer 量化 v1 thrashing 实际代价 | 5 min |
+| 3 | 实现 v2-fix-1（reset-on-success）| 30 min |
+| 4 | smoke test | 10 min |
+| 5 | 完整 v2 run（KVC 1P3D ts=1 N=1）| ~5.5h |
+| 6 | 三方对比：baseline / v1 / v2 | 30 min |
+| 7 | 决定是否需要 v2-fix-2 / v2-fix-3 | – |
+
+---
+
+## 8. 三方对比预测（待数据验证）
+
+| 指标 | baseline（v0）| v1（thrashing）| **v2（self-healing 预测）** |
+|---|---:|---:|---:|
+| Errors | 5 | ? | 2-5（仅 35680/39360 等真容量超限）|
+| Per-D 均衡 | ±26% | **±1.1%** | ±5-10%（部分 pin 仍 sticky）|
+| Direct-to-D rate | 42.8% | ?（可能因 thrash 反而下降）| **65-75%**（持续 affinity，转换 §1 fallback）|
+| Lat mean | 1.574s | ?（可能因 thrash 上升）| **1.30-1.45s**（达到 4DP 1.443s 水平）|
+| TTFT mean | 0.244s | ? | **0.10-0.15s** |
+| 最大 D-switches/session | 0 | 116 | <10（仅真饿死 session）|
+| Sessions 永久饿死 | 18 | 0 | 2-3（仅真容量超限）|
+
+预测核心：v2 应该结合 baseline 的稳定性（70-turn streak 应保留）+ v1 的公平性（无永久饿死），消除 v1 的 thrashing 副作用。
+
+---
+
+## 9. 局限与未验证
+
+1. **v1 中期数据 (23%) 推测**：完整数据可能改变 thrashing 严重性的判断
+2. **session 6880 trajectory 的崩溃机理是推断**：基于 admission events 数据 + streak 模式，但没有直接日志证明 reject 计数何时跨阈值（需要在 v2 加 instrument 输出）
+3. **reset-on-success 的预测效果未验证**：基于"70 turn 成功" + "1-2 次瞬时 reject" 的假设；如果 burst 持续多 turn，仍可能跨阈值
+4. **可能还有未发现的设计 bug**：v2 也许还会暴露新问题
+5. **三方对比需 same trace + same scale + same ts=1**：baseline 已有 N=3，v1/v2 各 N=1（ts=1 确定性 → N=1 可信）
+
+---
+
+## 10. 给 TEAM_REPORT 和 REFACTOR_PLAN_V1 的更新建议
+
+完成 v2 验证后：
+
+1. 在 `TEAM_REPORT` §3 ts=1 验证更新章节加入 §3.3 "Migration mechanism evolution: v0 → v1 → v2"
+2. 在 `REFACTOR_PLAN_V1` §6.2 标注实施反思——预设的 "rejection blacklist" 设计漏掉了 reset-on-success 这条
+3. 在新文档 `docs/POLICY_DESIGN_PRINCIPLES_ZH.md` 提炼出原则："任何会累积的代价机制必须配 healing/decay 机制，否则会陷入 self-amplifying 失效模式"
+
+---
+
+## 附录 A：本文数据来源
+
+| 章节 | 数据源 |
+|---|---|
+| §2 | `outputs/qwen3-30b-tp1-ts1-migration-v1/kvcache-centric-*/` 中期日志 |
+| §3.1 | `structural/session-d-binding.jsonl` 跨 turn 序列 |
+| §3.3 | `structural/admission-events.jsonl` mode/reason 交叉表 |
+
+## 附录 B：相关代码位置
+
+| 内容 | 位置 |
+|---|---|
+| RoutingState.session_d_rejects | `src/agentic_pd_hybrid/policies.py:46` |
+| KvAwarePolicy.select 跳过 blacklisted D | `src/agentic_pd_hybrid/policies.py:155-162` |
+| Degenerate fallback 选最少拒的 D | `src/agentic_pd_hybrid/policies.py:184-192` |
+| record_admission_reject 触发位置 | `src/agentic_pd_hybrid/replay.py:359-364`（_run_request） |
+| _is_admission_rejection_mode 子串集合 | `src/agentic_pd_hybrid/replay.py` `_ADMISSION_REJECTION_SUBSTRINGS` |
+| _fallthrough_reason 分类 | `src/agentic_pd_hybrid/replay.py` `_fallthrough_reason` |

docs/PROJECT_OVERVIEW.md

@@ -0,0 +1,98 @@
+# 项目概览
+
+这个项目验证一个问题：
+
+**agentic coding workload 里，如果 router 更懂 session 和 KV cache，P/D serving 的端到端延迟能不能更低。**
+
+当前基于：
+
+- SGLang `v0.5.10`
+- Qwen3-Coder-30B-A3B-Instruct
+- 单机 8 GPU
+- Mooncake loopback 模拟 P -> D 传输
+
+## 设计
+
+代码按两层分开：
+
+- **机制**：启动 SGLang、发送请求、管理 session、收集 metrics。
+- **策略**：决定请求去哪个 P node、哪个 D node。
+
+这样后续可以单独改 routing policy，不把它和 SGLang/xPyD 机制混在一起。
+
+## 已实现
+
+- 单机 P/D stack 启动和关闭。
+- 本地 Python PD router。
+- Ali trace 加载、session 级采样、synthetic prompt 生成。
+- 按 trace 原始到达时间 replay，不用固定 concurrency 强行压流量。
+- request-level metrics 和 summary。
+- 路由策略：
+  - `default`
+  - `sticky`
+  - `kv-aware`
+- serving 机制：
+  - `pd-disaggregation`
+  - `kvcache-centric`
+  - `pd-colo`
+- micro-benchmark trace 生成。
+- worker-managed / router-managed KV admission 对比。
+- worker-managed 下的 D session soft-cap，避免所有 session 都挤进 D KV。
+- SGLang patch：
+  - decode worker 支持 PD mode 下 local append-prefill；
+  - 暴露 streaming session cache 状态；
+  - 支持按 session 粒度 evict idle streaming session；
+  - 支持 direct append admission 查询。
+
+## 当前结论
+
+micro-benchmark 上，`kvcache-centric` 可以比 `pd-disaggregation` 好。
+
+原因很简单：session 少，D KV 放得下，turn2+ 可以直接走 D session，省掉一部分 P/D 路径开销。
+
+但在 300+ request、58 session 的测试上，情况不同：
+
+- D KV 放不下全部 session working set。
+- naive worker-managed 会频繁 evict/reseed 整个 session。
+- reseed 和 transfer 压力会抵消 KV reuse 收益。
+- aggressive P-backup 会增加尾延迟风险。
+
+当前 soft-cap 优化后：
+
+- worker-managed 比旧版本更稳；
+- TTFT 明显下降；
+- 没有再出现 600s transfer hang 被当成成功响应的问题；
+- 但 sampled Ali trace 上，`pd-disaggregation` 仍然略好。
+
+当前判断：
+
+**KV-cache-centric 只应该保留真正 hot 的 session。不是所有 session 都值得占 D KV。**
+
+下一步最有价值的是：
+
+- inter-turn-gap-aware admission；
+- session aging；
+- 更精确地预测哪些 session 会很快复用 KV。
+
+## SGLang 维护方式
+
+`third_party/sglang` 已纳入主仓库。
+
+历史结构：
+
+- `chore: vendor sglang v0.5.10 snapshot`：干净上游基线。
+- `feat(sglang): ...` / `fix(sglang): ...`：我们的 SGLang patch。
+
+后续改 SGLang 时：
+
+- 只改 `third_party/sglang` 下相关文件；
+- 单独提交；
+- commit message 带 `(sglang)`；
+- 不把 benchmark 输出、pyc、日志混进提交。
+
+## 已知限制
+
+- 这是实验原型，不是生产 router。
+- 当前主要验证单机 8 GPU。
+- Ali trace 没有原始 prompt，只能用 `hash_ids` 合成 prompt。
+- 当前 routing 还缺少真正的 hot-session 预测。

docs/REFACTOR_PLAN_V1_ZH.md

@@ -0,0 +1,385 @@
+# Refactor Plan v1：基于 ts=1 验证后的重构方向
+
+**日期**：2026-05-08
+**前置文档**：
+- `docs/REFACTOR_PLAN_ZH.md`（v0，已被本文 supersede——v0 的 backpressure 切入点结论已撤回）
+- `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md`（包含 §1-§7 结构性问题清单）
+- `docs/STRUCTURAL_VALIDATION_REPORT_ZH.md`（ts=10 数据下的早期验证）
+
+**触发**：`outputs/qwen3-30b-tp1-ts1-validation/` 4 个 run 完成（KVC 1P3D × N=3 + 4DP CA × 1，全部 ts=1）
+
+**目的**：把 ts=1 验证结果落到具体的重构决策——哪些事必须做、哪些事不要再做、KVC 项目本身是否需要重新定义价值主张
+
+---
+
+## 0. TL;DR
+
+1. **ts=10 失真是真的，影响 5-10×**——KVC 在 ts=10 灾难性输 DP 是 benchmark artifact，不是机制本身有问题
+2. **ts=1 同 scale 下 KVC ≈ DP**：lat mean 差 9%，TTFT 差 47%，errors 双 0
+3. **TEAM_REPORT 的 §1（session pin 不公平）是真问题，但代价从 6× 降到 ~2×**——仍是唯一值得做的 KVC 优化
+4. **TEAM_REPORT 的 §2/§3/§4/§5 大多是 ts=10 高压 artifact**——ts=1 下要么不显著、要么自然吸收
+5. **N=1 不可信是 ts=10 现象**——ts=1 下系统在 categorical 层面完全确定（routing/admission/errors 三次 run 完全相同）
+
+**项目落到情景 B（KVC ≈ DP）**——三种 forward 路径任团队决策（见 §6）。
+
+---
+
+## 1. ts=1 验证数据
+
+### 1.1 实验配置
+
+| 项 | 值 |
+|---|---|
+| Trace | `outputs/qwen35-swebench-50sess.jsonl`（4449 reqs / 52 sessions） |
+| 模型 | Qwen3-30B-A3B-Instruct-2507（TP1） |
+| 硬件 | 单机 4× H100 80GB（注：原始 ts=10 实验是 8 GPU；本次缩配） |
+| Time-scale | 1（真实 trace 时序，inter-turn gap p50 = 2.5s） |
+| Concurrency | 32 |
+| KVC 配置 | 1P3D，policy=kv-aware，admission=worker，seed-min-turn=1，prefill-priority-eviction |
+| DP 配置 | 4-way colo，policy=kv-aware（cache-aware） |
+| 输出根 | `outputs/qwen3-30b-tp1-ts1-validation/` |
+
+### 1.2 Headline 对比
+
+| Metric | KVC 1P3D ts=1（N=3 均值）| 4DP ts=1 | Delta |
+|---|---:|---:|---:|
+| **真实 mechanism errors** | **0** | **0** | 平 |
+| 报告 errors（口径不一致，见 §1.3） | 5 | 0 | – |
+| Lat mean | 1.574s | **1.443s** | DP 优 9% |
+| Lat p50 | 0.810s | **0.659s** | DP 优 19% |
+| Lat p90 | 3.796s | **3.641s** | DP 优 4% |
+| Lat p99 | 8.722s | **8.433s** | DP 优 3% |
+| TTFT mean | 0.244s | **0.129s** | DP 优 47% |
+| TTFT p50 | 0.122s | **0.090s** | DP 优 26% |
+| TTFT p90 | 0.572s | **0.252s** | DP 优 56% |
+| Per-worker spread | ±3.8% (3D) | ±3.1% (4 direct) | 接近 |
+
+### 1.3 KVC 5 errors 的真实身份
+
+DP 的同 5 个 (sess, turn) 也"失败"——但 metrics 口径不同：
+
+```
+KVC: 计入 error_count
+DP:  metrics 记 error=OK + finish_reason={'type':'abort', 'message':'Input length (X) exceeds the maximum allowed length (87811)'}
+```
+
+| sess | turn | input_len | KVC max | DP max |
+|---|---:|---:|---:|---:|
+| 35680 | 132 | 91600 | 92098 (✓) | 87811 (✗) |
+| 35680 | 133 | 92335 | 92098 (✗) | 87811 (✗) |
+| 39360 | 137 | 91700 | 92098 (✓) | 87811 (✗) |
+| 39360 | 138 | 92003 | 92098 (✓) | 87811 (✗) |
+| 39360 | 139 | 92135 | 92098 (✗) | 87811 (✗) |
+
+**两边都拒同样的请求**——区别只在于 KVC 在 P 端拒（KV 池满）、DP 在 prefill 端拒（max-input limit）。**真实 mechanism 错误率：KVC 0 / DP 0**。
+
+### 1.4 ts=1 的确定性
+
+KVC N=3 三次 run 跨 4449 records：
+
+| 维度 | 跨 run 差异 |
+|---|---|
+| `execution_mode` | **0 / 4449** records 不同 |
+| `assigned_decode_node` | **0 / 4449** records 不同 |
+| Errors（5 个 sess/turn 对） | **完全相同** |
+| 18 starved + 16 lucky session | **完全相同** |
+| Per-D load (1502/1445/1502) | **完全相同** |
+| Lat mean | 1.574 / 1.573 / 1.574（**0.06%** 漂移）|
+| Lat p50 | 0.811 / 0.809 / 0.812（**0.4%** 漂移）|
+| 单 request lat | abs p90 diff = 25ms |
+
+**结论**：低压 / ts=1 区间下 KVC 系统在 categorical 层面（路由 / admission / 失败位置）**完全确定**，仅低层数值有 model 计算微抖动。
+
+---
+
+## 2. 对 TEAM_REPORT §1-§7 的修订
+
+| § | TEAM_REPORT 原 claim | TEAM_REPORT 原优先级 | ts=1 验证后状态 | **修订优先级** |
+|---|---|---|---|---|
+| §2.1 | session pin + 容量盲选 → 25% 饿死 | **P0** | ✅ 结构性问题仍在（18/52 session 永久 pin），但代价从 6× 慢降到 ~2× | **P0**（唯一值得做的 KVC 优化）|
+| §2.2 | D-side LRU 跟不上 → 8% errors | **P0** | ⚠️ D 仍瞬时顶到 token_usage=1.00，但**ts=1 下 drain time 自然吸收**——0 KVTransferError 雪崩（vs ts=10 369 次） | **降级 P3**（drain time 已解决症状）|
+| §2.3 | 无 backpressure 通道 | P1（已实现）| ❌ ts=1 下 transfer cascade 不存在，backpressure 无作用对象 | **冷藏**（代码留着，但默认 off）|
+| §2.4 | P-side round-robin 不感知 D 健康 → prefill-0/-1 错误差 180× | P1 | ⚠️ 1P 配置不可测；ts=10 现象**高度怀疑也是 artifact**（错误本身在 ts=1 消失） | **存疑 / 重测后再说** |
+| §2.5 | admission RPC 进 scheduler 主循环 → 1Hz polling 让 errors ↑46× | P2 | ❌ 是 ts=10 高压时的现象，ts=1 下不显著 | **冷藏** |
+| §2.6 | time-scale=10 失真 → 所有 KVC vs DP 结论可能被放大 | **P0** | ✅ **完全证实**（74× errors↓, 8.7× TTFT↓, 7× per-D spread↓） | **DONE，作为前置条件锁定** |
+| §2.7 | execution_mode 标签命名错位 | P1 | ✅ 仍存在；本次 ts=1 又发现 `error_count` 在 KVC vs DP 口径不一致 | **P1**（纯 labeling 修复，~半天）|
+| §2.8 | N=1 不可信 → 实验必 N≥3 | P2 | ⚠️ **是 ts=10 高压现象**——ts=1 下 N=1 categorical 完全确定 | **改写规则**：高压 N≥3 / 常规 N=1 |
+| §2.9 | microbench 把 KVC 失效条件全规避 | – | 仍成立 | **保留观察**（实验设计原则）|
+
+---
+
+## 3. v0 REFACTOR_PLAN 回顾
+
+### 3.1 v0 做对的
+
+- **唯一代码改动选 backpressure**：作为对 §2.3 的最小验证手段是合理的
+- **预算 KISS**：用 8h GPU 验证 §1-§7，思路正确
+- **明确"P0 是 time-scale=1 baseline"**：v0 的 §1 末尾就指出 "time-scale=1 验证为 P0 待办"——本次实验正是把这条做了
+
+### 3.2 v0 的核心误判
+
+| v0 假设 | 实际 |
+|---|---|
+| backpressure 是 §3 的最小验证 → 也是修复 | ts=1 下 §3 的症状（transfer cascade）不存在，backpressure 无效 |
+| 8h 预算够跑 ts=1 baseline + backpressure smoke | ts=1 单 run 5.5h，4 run 全跑要 22h（实际跑了 22h） |
+| §1 / §2 的修复"超出 KISS 边界"，先验证不修 | 验证后发现 §1 是**唯一**值得做的真问题，应该早点把它纳入 |
+
+### 3.3 v0 的 backpressure 代码命运
+
+代码保留（`--enable-backpressure` 默认 off），原因：
+- 不删除是因为如果未来跑高压 / 大 trace / 真 RDMA 失败回归到类 ts=10 区间，可能仍有用
+- 但**不部署、不启用、不文档化为推荐配置**——避免给以后看到代码的人误导
+
+---
+
+## 4. 修订后的优先级矩阵
+
+```
+                    必做                   建议做                  不做
+                  ────────              ────────              ────────
+ts=1 必修        §1 capacity-aware   (空)                   §2 / §3 / §4 / §5
+                 policy + migration                          的 ts=10 fix
+
+ts=1 nice        §2.7 metrics 标签   (空)                   §2.8 N≥3 严苛规则
+to have         统一口径                                    （改成"高压 N≥3"）
+
+文档              §3 写入 TEAM      v0 标记 superseded     ts=10 数据归档
+                  REPORT 更新                               （但保留可追溯性）
+```
+
+**唯一进入"必做工程"列表的是 §1**。其他全是文档或冷藏。
+
+---
+
+## 5. KVC vs DP 拆分到 path-level 看真实差距
+
+理解 §1 的 ROI 必须先看 path-level（不是整体均值）：
+
+### 5.1 KVC 内部 path 性能（来自 ts=1 N=3 一致数据）
+
+| Path | n | 占比 | Lat p50 | TTFT p50 |
+|---|---:|---:|---:|---:|
+| `kvcache-direct-to-d-session`（快路径）| 1903 | **42.8%** | **0.475s** | **0.042s** |
+| `pd-router-fallback-large-append-session-cap`（慢路径）| 2409 | **54.2%** | 1.04s | 0.32s |
+| `pd-router-turn1-seed`（每 session 一次）| 52 | 1.2% | 0.375s | 0.057s |
+| 其余 | 85 | 1.8% | 多种 | 多种 |
+
+### 5.2 DP 全部 path（单一）
+
+| Path | n | 占比 | Lat p50 | TTFT p50 |
+|---|---:|---:|---:|---:|
+| `dp-colo-router` | 4449 | 100% | 0.659s | **0.090s** |
+
+### 5.3 路径级对比
+
+| | KVC direct | KVC fallback | DP |
+|---|---|---|---|
+| Lat p50 | **0.475s**（赢 DP 28%）| 1.04s（输 DP 58%）| 0.659s |
+| TTFT p50 | **0.042s**（赢 DP 53%）| 0.317s（输 DP 252%）| 0.090s |
+
+**事实陈述**：
+- KVC 快路径 **明显快于** DP（无 P 介入、无 mooncake transfer）
+- KVC 慢路径 **明显慢于** DP（P→D transfer 开销没法摊到 turn 内）
+- 当前 quick:slow = 42.8% : 54.2%——慢路径多 → 整体输 DP 9-47%
+- 如果能把比例反过来到 70:25 或更好，KVC 整体会赢 DP
+
+**§1 的本质就是"为什么有 54% 进了慢路径"**——因为 18/52 session 被 pin 在容量紧张的 D 上，每次 admission 都拒。
+
+---
+
+## 6. 三种 forward 路径
+
+> **更新（2026-05-09）**：情景 C **已实现**——见 `docs/V2_RESULTS_ZH.md`。下面三个分支保留作历史记录。
+>
+> | 情景 | 描述 | 状态 |
+> |---|---|---|
+> | A | KVC < DP，接受现状转维护 | 不适用 |
+> | B | KVC ≈ DP，重新定义价值主张 | 不适用 |
+> | **C** | **KVC > DP，优化拉大差距** | **✓ 实现：v2 在 7/8 头部指标击败 4DP（TTFT mean -24%, p50 -54%, p90 -64%；lat mean -0.8%, p50 -12.6%）** |
+>
+> 关键修复：(1) reset-on-success blacklist decay（消除 v1 thrashing），(2) `--kvcache-direct-max-uncached-tokens` 2048→8192（让 41% 大 append 走 direct-to-D 快路径）。direct-to-D rate 从 baseline 42.8% 升到 v2 91.7%。
+
+### 6.1 选项 A：接受现状，项目转维护
+
+**判断**：KVC 在 ts=1 + 同 scale 下 ≈ DP（9% 慢、47% TTFT 慢），但**也没灾难性输**。如果项目目标是"验证 KV-aware routing 在 agentic 上是否可行"，答案是 **可行但收益不显著**。
+
+**操作**：
+- 写 TEAM_REPORT §3 总结 ts=1 实验
+- 把 ts=1 数据 + 4 个 run 归档到 `RESULTS_FROZEN_TS1.md`
+- KVC 代码保留但标记 "experimental, not recommended for production"
+- 团队转下一个项目方向（不是本文范围）
+
+**成本**：1 周文档收尾。
+**风险**：放弃了 §1 修复后可能的 KVC > DP 上限。
+
+### 6.2 选项 B：做 §1，目标让 KVC > DP
+
+**判断**：5.3 节的路径分析表明 KVC 快路径已经赢 DP；如果把饿死 session 救回快路径，KVC 整体可能赢 DP。
+
+**具体改动**：
+
+#### 6.2.1 capacity-aware policy（`policies.py:166-172`）
+
+当前评分（无容量项）：
+```python
+score = (
+    overlap + sticky * self.sticky_bonus,
+    sticky,
+    inflight_penalty,
+    assignment_penalty,
+)
+```
+
+提议改为：
+```python
+# 新增：D 当前容量利用率（从 worker-mode admission 已能查到）
+capacity_used = worker_capacity_used_ratio.get(worker.worker_id, 0.0)
+
+# Hard cap：容量 > X 时禁止该 D 进入候选
+if capacity_used > HARD_CAP_THRESHOLD:  # e.g. 0.85
+    continue
+
+score = (
+    overlap_capped,           # 原 overlap，但限幅避免单个 D 永远赢
+    -capacity_used,           # 新增二级排序项：偏好空闲 D
+    sticky,
+    inflight_penalty,
+)
+```
+
+#### 6.2.2 session migration（`replay.py` 或 policy 层）
+
+当 session X 在 D-A 上连续被 admission 拒 N 次（如 N=3）：
+- 主动 release X 在 D-A 上的 session state
+- 允许下次 turn 把 X 路由到另一个 D
+- 代价：丢失 D-A 上已积累的 KV——但 fallback 路径本来也丢了，**净收益正**
+
+#### 6.2.3 metric 修复（`replay.py`）
+
+把"`pd-router-fallback-large-append-*`" 标签按真实原因细分：
+- `session-not-resident-on-pinned-D`（§1 主因）
+- `real-large-append`（>2048 阈值，§2.7）
+- `session-was-evicted`（被 LRU 踢过）
+- `session-cap-rejected`（worker admission 拒）
+
+让以后看 metrics 的人不再被名字误导。
+
+#### 6.2.4 验证
+
+- 每改动跑 KVC 1P3D ts=1 N=1（categorical 确定，不需要 N=3）
+- 对比 baseline run1（已有数据）
+- 关键指标：`kvcache-direct-to-d-session` 占比、整体 lat mean、TTFT mean
+- 目标：direct-to-D rate 从 42.8% 升到 > 70%、整体 lat 追平或赢 DP
+
+**成本**：3 天编码 + 5 天测试 + 2 天文档 ≈ 2 周。
+**风险**：
+- session migration 可能导致 thrash（A→B→A→B），需要冷却时间机制
+- capacity HARD_CAP 阈值需要 sweep 找最优
+- 改完仍可能不赢 DP（理论上限不知道）
+
+### 6.3 选项 C：保留 KVC，但寻找 KVC 真正赢的工作点
+
+**判断**：当前 SWE-Bench 50 sessions × 30B 模型 × 4 GPU 是一个特定工作点。KVC 的设计初衷是"长 multi-turn session 的 KV 复用"——可能在某些其他工作点有显著优势。
+
+**候选工作点**：
+- **更长 session（>200 turns）**：复用收益更大
+- **更小模型（如 7B / 14B）**：mooncake transfer 占比更大，KVC 节省更明显
+- **更大 trace（>200 sessions）**：DP 的 prefix cache 命中率会下降，KVC 的 session-aware 优势放大
+- **真实 RDMA（非 mooncake TCP loopback）**：transfer 更快，KVC 的 P→D 开销更小
+
+**操作**：
+- 设计 1-2 个新 micro/macro benchmark
+- 跑 KVC vs DP 对比
+- 找到差距 > 30% 的工作点（KVC 赢 / 输都是数据）
+
+**成本**：~1 个月（trace 设计 + benchmark + 分析）。
+**风险**：可能找不到 KVC 显著赢的工作点。
+
+---
+
+## 7. 推荐组合
+
+按风险 / 收益排序：
+
+1. **必做**（无论选 A/B/C）：
+   - 写 `TEAM_REPORT §3 ts=1 验证更新`
+   - 修 `metrics 标签口径`（§2.7 + KVC/DP error_count 一致化）
+   - **冷藏 backpressure 代码**（不删但默认 off）
+   - 把 v0 REFACTOR_PLAN 标 superseded
+
+2. **强烈推荐**：选项 B 的 §6.2.1（capacity-aware policy hard cap）
+   - 工程量小（~1 天编码 + 1 天测试）
+   - 验证 §1 修复的真实收益是否如预测
+   - 如果 direct-to-D rate 不显著提升 → 把 §6.2.2 也加上
+   - 如果还不行 → 接受现状走选项 A
+
+3. **看团队带宽**：选项 C 的工作点探索
+   - 不与 §6.2 冲突，可以并行
+   - 找到一个 KVC 真正赢的工作点会极大改变项目价值主张
+
+---
+
+## 8. 应该砍掉的事（明确列表）
+
+| 事 | 砍的理由 |
+|---|---|
+| backpressure smoke sweep（v0 计划的 4 run） | ts=1 下 backpressure 无作用对象 |
+| §2.5 admission API probe/commit 拆分 | 高压才显著，等找到 KVC 高压 workload 再说 |
+| §2.2 D-side 分层 LRU eviction（hot retract） | drain time 自然吸收 |
+| §2.4 P-side D-health-aware routing | 1P 测不出，ts=10 现象高度存疑 |
+| 大量 instrument（admission-events / pool timeseries） | 已经够了，先用现有数据 |
+| 任何 ts=10 区间的优化 | 那是 benchmark artifact 主导的区间，不代表真实部署 |
+| N≥3 实验作为硬规则 | 改写为"高压 N≥3，常规 N=1 即可" |
+
+---
+
+## 9. 风险与未验证的假设
+
+1. **4DP ts=1 是 N=1**：虽然 KVC ts=1 是确定性的，DP 是新机制 N=1，理论上需要 N≥3 验证。但 DP 在 ts=10 也是 0 errors / 1.43s mean，行为相对 KVC 更稳定，N=1 风险较小。**如选项 B 推进，建议补 N=2**。
+2. **2 个 input-too-long session 是 trace 数据问题**：这两个 session（35680、39360）在 turn 132+ / 137+ 才超过 input limit。可能是 trace 生成时没控制好 max input。**应该独立把这两个 session 从 trace 移除或截断后重跑作为对照**。
+3. **4 GPU 缩配 vs 8 GPU 原始**：本次 1P3D / 4DP 数据无法跨 8 GPU 原始数据直接比，需要在结论中明确。但 ts=1 + 同 scale 内部对比是干净的。
+4. **mooncake TCP loopback**：所有 transfer 在单机 TCP 模拟下进行。生产 RDMA 下 KVC 的 transfer 开销可能显著降低，KVC 优势可能扩大——这是 **选项 C 的一个候选维度**。
+5. **§1 修复是否真能让 direct-to-D 上升到 70%+ 是预测**：实际可能受 hash overlap 限制（即使 D 容量充裕，没有 prefix overlap 就走不了 direct-to-D）。**需要 §6.2 验证后才知道天花板**。
+6. **input-limit error 的 metrics 口径修复影响以后所有比较**：注意修改后 ts=10 历史数据的 error_count 也需要重算（或在分析时显式补偿）。
+
+---
+
+## 10. 决策点（需要团队确认）
+
+请审阅后回答：
+
+| # | 决策 | 选项 |
+|---|---|---|
+| D1 | 选哪条 forward 路径？ | A（维护）/ B（修 §1）/ C（探索 workload）/ B+C |
+| D2 | 写 TEAM_REPORT §3 ts=1 验证更新章节？ | Yes / No |
+| D3 | 把 v0 REFACTOR_PLAN 标 superseded？ | Yes / No |
+| D4 | 删除 backpressure 代码 vs 冷藏？ | 删 / 冷藏（默认 off）|
+| D5 | 修 metrics 标签口径（§2.7 + error_count 一致化）？ | Yes / No |
+| D6 | 是否补 4DP ts=1 N=2 / N=3 做更稳的 baseline？ | Yes / No |
+| D7 | 是否把 sess 35680 / 39360 从 trace 移除做"干净" baseline？ | Yes / No |
+
+---
+
+## 附录 A：本文数据来源
+
+| 章节 | 数据源 |
+|---|---|
+| §1.2-§1.4 | `outputs/qwen3-30b-tp1-ts1-validation/{kvc_1p3d_run{1,2,3},dp4}_{summary.json,metrics.jsonl}` |
+| §1.4 跨 run 一致性 | per-record diff via `scripts/analysis/analyze_ts1_validation.py` + 临时 diff 脚本 |
+| §5 path-level | metrics.jsonl 按 `execution_mode` 分组 |
+| §2 §1-§7 修订 | `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md` 原数据 + ts=1 新数据交叉对比 |
+
+## 附录 B：相关文档
+
+- `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md` — §1-§7 原结构性问题清单
+- `docs/REFACTOR_PLAN_ZH.md` — v0 重构计划（本文 supersede）
+- `docs/AGENTIC_FIT_ANALYSIS_ZH.md` — 早期 fit 分析（§1-§7 来源）
+- `docs/STRUCTURAL_VALIDATION_REPORT_ZH.md` — ts=10 结构性 claim 验证
+- `docs/KVC_DEBUG_JOURNEY_V1_TO_V5.md` — v1→v5 演进
+- `docs/V5_PROFILE_INVESTIGATION_ZH.md` — v5+profile 调查（已 critic 修订）
+- `scripts/sweep_ts1_kvc_n3_plus_dp.sh` — 本次 4 run sweep 脚本
+- `scripts/analysis/analyze_ts1_validation.py` — 本次分析脚本
+
+---
+
+**作者注**：本文偏决策导向。如果要写更技术的 §1 capacity-aware policy 实现细节，应该在 D1 决策为 B 之后单独出一份 `IMPL_CAPACITY_AWARE_POLICY.md`。

docs/REFACTOR_PLAN_ZH.md

@@ -0,0 +1,123 @@
+# Refactor Plan v0：极简版
+
+**日期**：2026-05-06
+**目标**：用最小改动 + 轻量实验，验证 `docs/AGENTIC_FIT_ANALYSIS_ZH.md` 提出的结构性缺陷是否真实存在、影响多大。
+**预算**：8h GPU 时间（约 4-6 次 ~30-60 min smoke run）。
+**KISS 边界**：不动 SGLang `scheduler.py` 主循环结构；不引入新 mooncake 协议；不实现 cross-D session migration；不做 admission probe/commit 拆分；不动 LRU eviction 策略。
+
+## 计划结论（与用户已确认的）
+
+回审 plan-v0 时发现两个原 Phase 1 改动**都不是 bug**：
+
+- `_estimate_session_resident_tokens` 返回 full prompt 是设计如此——所有需要"增量"的 call site 都已经做 `target - current` 减法（`replay.py:1247-1254`、`:1393-1394`、`:1490-1491`）。
+- `decode_resident_blocks` 不缩减只是浪费几 MB 内存，**不影响 routing 决策**（SWE trace 的 hash_ids 是 session-unique，policy 仍能正确选 D）。
+
+最终极简版只做一件代码改动（**加 backpressure**）+ 大量 instrumentation。
+
+## 唯一代码改动：Backpressure 信号
+
+### 改动点 1：SGLang `admit_direct_append` 响应增加两个字段
+
+文件：`third_party/sglang/python/sglang/srt/managers/io_struct.py`、`scheduler.py`
+
+```python
+@dataclass
+class DirectAppendAdmissionReqOutput:
+    ...                         # 已有字段保留
+    recommended_pause_ms: int = 0   # 新增
+    queue_depth: int = 0            # 新增
+```
+
+`scheduler.py:admit_direct_append` 末尾计算 hint：
+
+```python
+def _compute_backpressure_pause_hint(self) -> float:
+    depth = len(self.disagg_decode_transfer_queue.queue)
+    if depth < 8:
+        return 0.0
+    return min(2000.0, depth * 100.0)   # 简单线性
+```
+
+### 改动点 2：replay 端按 hint 退避
+
+文件：`src/agentic_pd_hybrid/replay.py`
+
+- `DecodeResidencyState` 新增 `pause_until_s: dict[str, float]`
+- `_query_decode_direct_admission` 解析响应里的 `recommended_pause_ms`，更新 `pause_until_s[server_url] = now + pause_ms / 1000`
+- 在调 `_invoke_router` / `_invoke_decode_session_direct` 前检查 `pause_until_s[decode_url]`，若 `now < pause_until` 则 sleep 到该时刻
+
+### 改动点 3：新 CLI flag
+
+`src/agentic_pd_hybrid/cli.py`、`benchmark.py`：
+
+```
+--enable-backpressure   # 默认 false，保留 baseline 行为
+```
+
+### 改动点 4：观测日志
+
+每个 run dir 新增三个 jsonl：
+
+- `admission-events.jsonl`：每次 admission RPC（timestamp, session, D, can_admit, queue_depth, pause_ms, latency_s, available_tokens, evicted_session_count）
+- `backpressure-events.jsonl`：每次实际 sleep（timestamp, D, sleep_ms, queue_depth_at_signal）
+- `session-d-binding.jsonl`：每个 session 第一次 open 在某 D 时记录（timestamp, session, D, turn_id）
+
+## 实验矩阵（8h 预算内）
+
+按"先做 anchor，再做单变量对照"排序。每行右侧是预估机时。
+
+| ID | 配置 | 目的 | 机时 |
+|---|---|---|---|
+| **E0 (existing)** | v5 baseline，time-scale=10，无 backpressure | Anchor，已存在 `outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/run1` | 0 |
+| **E1** | v5 + backpressure ON，time-scale=10，全 trace | 验证 Claim §3（backpressure 是否能消除 KVTransferError 雪崩） | ~50 min |
+| **E2** | v5 baseline，time-scale=1，**短 trace**（前 12 sessions ≈ 1000 reqs） | 验证 Claim §7（time-scale=10 失真）；不开 backpressure | ~60 min |
+| **E3** | 8DP CA，time-scale=1，同 E2 trace | E2 的对照——真实时序下 KVC 是否仍输 DP | ~60 min |
+| **E4** | v5 + backpressure，time-scale=1，同 E2 trace | backpressure 在真实时序下还有用吗？ | ~60 min |
+| **E5**（备选） | v5 baseline，time-scale=10，**concurrency=4**，全 trace | 验证 Claim §1（高并发是不是必要条件） | ~50 min |
+
+总：4-5 个 run，~3-5h。剩余预算给失败重跑/分析。
+
+## 实验目标——回到 §1-§7 一一对照
+
+| 文档 § | Claim | 由哪个 exp 证伪/支持 | 需要的指标 |
+|---|---|---|---|
+| §1 | Session 永久 pin + 容量盲选造成双峰 | 已有 E0 数据足够 | direct-to-D rate per session distribution |
+| §2 | LRU 跟不上压力 | 已有 E0 logs 足够 + E1 看 backpressure 之后 trim/error 比例变化 | trim 事件数 vs OOM 数 |
+| §3 | 没 backpressure 是雪崩源 | E0 vs E1 | KVTransferError 数、P99 latency |
+| §4 | admission RPC 干扰 scheduler | 不在本轮实验范围（需要 admission probe 拆分才能验，不做） | – |
+| §5 | P-side 不感知 D 健康 | 已有 E0 logs 足够（prefill-0 vs prefill-1 错误数） | per-P KVTransferError |
+| §6 | (已撤回) | – | – |
+| §7 | time-scale=10 失真 | E0 vs E2（同 KVC，不同 time-scale）；E2 vs E3（同 time-scale，KVC vs DP） | latency 分布、direct-to-D rate |
+
+## Final 实验报告交付
+
+跑完后输出 `docs/STRUCTURAL_VALIDATION_REPORT_ZH.md`，按 §1-§7 每条给出：
+
+- **Claim 字面**
+- **数据证据**（哪个 exp、哪个 metric）
+- **结论**：成立 / 部分成立 / 推翻
+- **影响量化**：数字差异
+- **不确定性**：N=1 风险、其他 confounder
+
+## 不做的事（KISS 边界）
+
+| 想做但不做 | 理由 |
+|---|---|
+| 跑 N=3 重复 | 8h 装不下；single-run 可看大方向 |
+| 全 sweep 参数 | 只调 time-scale 和 backpressure 一个 boolean |
+| 改 LRU eviction | 不在本轮范围 |
+| Cross-D migration | 不在本轮范围 |
+| Admission probe/commit 拆分 | 不在本轮范围 |
+| P-side D-health routing | 不在本轮范围 |
+| 修两个"非 bug"（estimate / aging） | 验证后非真实 bug |
+
+## 预期失败路径
+
+- **GPU 资源紧张**：smoke trace 进一步压缩（前 8 sessions / 600 reqs）
+- **time-scale=1 跑超 1.5h**：截断到 600s 内能完成的部分
+- **backpressure 配错**：先用 sleep_ms = depth * 100 简单线性；调不通就回滚到 0（无 backpressure）
+- **SGLang patch 编译错**：所有 patch 在 io_struct.py 和 scheduler.py 的少量行内，可单独 git restore
+
+---
+
+接下来：实现 → 跑 smoke → 写报告。

docs/RESEED_SLOW_PATH_AND_D_TO_P_GAP_ZH.md

@@ -0,0 +1,368 @@
+# Reseed 慢路径现状与 D→P KV 同步缺口
+
+**日期**：2026-05-11
+**对象**：项目团队 + 后续 paper reviewer
+**性质**：基线现状落盘 + future-work 缺口定位
+**前置文档**：
+- `docs/V2_DEEP_ANALYSIS_ZH.md` §3.2 §4.2（reseed 路径在 v2 数据中的表现）
+- `docs/KVC_ROUTER_ALGORITHM.md` §3 §9（算法形式化 + open questions）
+
+**目的**：把"v2 的 reseed slow path 为什么慢、能不能用现有机制治、还差什么"三个问题落盘成单一参考文档，让团队不必再口头反复对齐，让论文 future-work 章节有可引用的基础。
+
+---
+
+## 0. TL;DR
+
+1. KVC v2 在 SWE-Bench 测试中 8.3% 请求走非 direct-to-D 的 reseed/fallback 路径，**单次 reseed 实测 3-7s**（TTFT p99 = 1.28s 全部来自这条路径）。
+2. 启用真 RDMA（节点有 mlx5_0/_1 @ 200 Gb/s × 2 active）能把 reseed 的 transfer 段（~1.5-4s）压到 ~200-400ms，但**对 re-prefill 段（~1.5-3s）无效**。预期 reseed 总时间从 3-7s 降到 1.7-3.2s，TTFT p99 ~0.7s，**仍输 DP（0.43s）**。
+3. 真正消除 reseed 长尾必须实现 **D→P 增量 KV 同步**——让 P 端 backup 跟上 D 在 direct-to-D append 路径上累积的 KV，避免 reseed 时重新跑 prefill kernel。
+4. 经 Opus agent 独立 forensic 审查（commit `9ccd853`）+ 全分支 git 检索：**当前代码、vendored SGLang、mooncake 三层均无 D→P 实现**，作者也没有在其它分支偷偷开发——仓库总共只有 main（旧 baseline）+ kvc-debug-journey-v1-to-v4（本工作分支）两个分支，main 还落后我们 18 个 commit。
+5. `--kvcache-prefill-backup-policy capacity-backup` 这个 flag 看起来像 D→P 同步但**不是**——它的真实语义只是"reseed 完不关 P streaming session"，P 端 KV 仍是 seed-time 的**静态快照**，不随 direct-to-D append 而增长。
+6. 实现 D→P 增量同步的工程量 ~1-2 周，最难的不是网络层（mooncake 加 D-sender / P-receiver 角色 ~400 LOC），而是 **SGLang radix tree 改成允许从外部 worker 喂数据**——radix cache 当前假设单一生产者。
+
+---
+
+## 1. 团队成员的三个质疑（关键框架，paper 引用建议保留原话）
+
+这三条质疑出自 v2 完成后的对话审查，**直接戳穿了"启用 capacity-backup 就能消除 slow path"的一厢情愿**。每条都有代码层证据支持，**全部成立**。
+
+### 质疑一：P 节点的 pool 塞得下所有 backup 的 KV cache 吗？
+
+**回答：塞不下，max 同时 backup ~1-2 个大 session。**
+
+代码证据（`src/agentic_pd_hybrid/replay.py:1618-1620`）：
+
+```python
+max_backup_sessions = max(1, capacity_tokens // max(1, target_tokens * 2))
+max_backup_sessions = min(max_backup_sessions, 4)
+```
+
+按 SWE workload 实测代入：
+- P 池 `capacity_tokens` ≈ 92,104 tokens（SGLang 启动时按 mem_fraction_static 自动分配）
+- 典型 session peak input `target_tokens` ≈ 50,000-80,000 tokens
+- 计算：`92K // (50K × 2) = 0` → `max(1, 0) = 1`
+- → **P 最多同时 backup 1 个大 session**
+
+对照小 session：
+- target 20K：`92K // 40K = 2` → backup 上限 2 个
+- target 10K：`92K // 20K = 4` → backup 上限 4 个（达到代码硬上限）
+
+→ **capacity-backup 在真实 agentic 长 context workload 下只能救少数 session，不是全员保险。**
+
+### 质疑二：P 上的 backup 是陈旧快照——49K 的 append 内容根本没经过 P
+
+**回答：完全正确，这是 capacity-backup 设计上的致命缺陷。**
+
+**用户提供的反例场景**（已成为 paper 中描述 slow path 的标准例子）：
+
+```
+turn 0:   P 做 prefill 1K tokens → 经 mooncake 传到 D → P 留 1K backup
+turn 1-50: 全部走 direct-to-D，D 上做 append-prefill，KV 在 D 上从 1K 增长到 50K
+           ↑↑↑ 关键：这 49K 的 append 内容（tool 输出、user 消息、模型生成）
+              **从未流经 P 节点**。P 端 backup 锁在 1K 状态。
+turn 51:  D 出于某种原因（容量、迁移、显式驱逐）拒绝 → 触发 reseed
+          → 即使 P 上有 backup，也只是 turn-0 的 1K
+          → 实际需要 D 上重建的是 50K（当前完整 context）
+          → P 必须从 prompt 重新 prefill 49K 的差额
+          → capacity-backup 节省的 compute 仅 ~2%
+```
+
+**代码证据**（独立 Opus agent forensic 审查，commit `9ccd853`）：
+
+1. 唯一更新 `session.prefill_resident_tokens` 的函数是 `_commit_prefill_backup_residency`（`replay.py:1483`）
+2. 这个函数的唯一 caller 是 `_invoke_kvcache_seeded_router`（`replay.py:2208`）—— 即 seed/reseed 路径
+3. `_invoke_session_direct`（`replay.py:2719`，direct-to-D 路径）只更新 `session.opened` / `resident_tokens` / `last_trace_request`，**从不触碰任何 P 端字段**
+4. `_commit_prefill_backup_residency` 内部用 `_estimate_session_resident_tokens(request)` 取的是**完整 request 的预估**，不是 append delta——所以连 bookkeeping 层面都不假设有增量更新
+
+→ **`capacity-backup` 的真实语义只是"reseed 完之后跳过 `_close_prefill_session`"**（`replay.py:2221`），P 端 streaming session 保持 open 状态、KV 留在 P 的 radix tree 中。但**不存在任何机制让这份 KV 跟上 D 端的 append 增长**。
+
+### 质疑三：D 触发 reseed 后，本机旧 session 的 KV cache 是不是清空了？P 做完 re-prefill，KV 推到哪里？
+
+**回答：是的，旧 KV 直接 free 掉；P 重新 prefill 完之后推到 router 选的新 target D（可能同 D，可能换 D）。中间没有"先 dump 到 P 再清"的快捷方式。**
+
+#### D 端驱逐时的 KV 处理
+
+代码证据（`replay.py:_close_decode_session`，1539-1569 行；`session_aware_cache.py:release_session`，250-276 行）：
+
+```python
+# replay.py 端
+async def _close_decode_session(..., evicting_for_capacity=False):
+    if not session.opened:
+        return
+    await _close_streaming_session(...)         # 给 D 发关闭信号
+    # 从 D 的 resident bookkeeping 里删掉这个 session
+    session.opened = False
+    session.resident_tokens = 0
+    if evicting_for_capacity and not session.prefill_opened:
+        residency.decode_evictions_without_prefill_backup += 1
+
+# SGLang 端（session_aware_cache.py）
+def release_session(self, session_id):
+    # 解锁引用 + 直接 free KV slots
+    self.token_to_kv_pool_allocator.free(kv_indices)
+    # ↑ 没有序列化、没有外发、没有 D→P 通道
+```
+
+**D 驱逐 = 把 KV slot 直接归还给 token pool 分配器。完全没有任何 outbound 网络调用。**
+
+#### Reseed 时 P→D 的目标选择
+
+驱逐之后的 reseed 路径（`_invoke_kvcache_seeded_router`，`replay.py:2101`）走的是与 turn 0 完全一样的 P-mediated seeding：
+
+```
+1. KvAwarePolicy.select() 选择一个 target D'（可能是同一个 D，也可能因 migration 换 D）
+2. _invoke_kvcache_seeded_router 在 D' 上 open 一个 streaming session
+3. 给 P 发完整 prompt → SGLang pd-router 让 P 做完整 prefill
+4. P 的 prefill 完成后通过 mooncake 把 KV 一次性推到 D'
+5. D' 上接收完毕，session 重建完成；decode 继续
+```
+
+**所以 P 做完 re-prefill 的 KV 推到 KvAwarePolicy 选的 target D'**——可能是：
+- 同一个 D（驱逐后重新接受）
+- 另一个 D（如果 reject 计数累积触发 migration，详见 KVC_ROUTER_ALGORITHM §3.3）
+
+无论哪种，**旧 D 的旧 KV 在新 KV 到达之前就已经被 free**。没有 D→D 的直接迁移路径，没有"先 dump 到 P 再推回"的快捷路径。
+
+---
+
+## 2. Reseed 路径的完整 step-by-step 现状
+
+把上面三个质疑串成端到端流程，以下是 v2 当前 reseed 路径的**完整**操作序列。每一步都标注实测耗时与代码位置。
+
+### 触发条件
+
+下列任一发生时 router 走 reseed 路径（详见 `KVC_ROUTER_ALGORITHM.md §3.3`）：
+- D 端 `Admit()` 返回 `can_admit=False`，原因为 `no-d-capacity` / `session-not-resident` / 等
+- KvAwarePolicy.select 返回的 D 不再持有该 session（migration 触发）
+- v1/v2 的 reject counter 累积让所有 D 都被 blacklist（极少触发，由 reset-on-success 保护）
+
+### 端到端时间线
+
+```
+t=0      上游 agent 发出 turn N 请求（input ~50K，append ~2K）
+            ↓
+t=~5ms   Router 的 KvAwarePolicy.select() 选 target D'（O(|D|) Python 评分）
+            ↓
+t=~10ms  Router → D' 发 admit_direct_append RPC
+            ↓
+t=~30ms  D' 返回 can_admit=False, reason="session-not-resident"
+         或 "no-d-capacity"，Algorithm 3 bump rejects[s, D']++
+            ↓ （fallback chain 最多再试 ε-1 个 D，对应 ε ~30ms 总额）
+t=~100ms 所有 D 都被拒 / 选不到适合 D，路径退化到 seeded router
+            ↓
+t=~110ms Router 转 _invoke_kvcache_seeded_router
+            ↓
+t=~120ms [可选] capacity-backup policy 下：_reserve_prefill_backup_capacity()
+         检查 P 池容量，若不够先 LRU 驱逐别的 P backup session
+            ↓
+t=~150ms P 上 open streaming session（HTTP /session/open）
+            ↓
+t=~200ms 发完整 prompt 到 SGLang pd-router → 路由到 P
+            ↓
+t=~250ms  P 开始 prefill
+         ↓
+         ↓ ←←← 大头 1：P-side re-prefill 段
+         ↓     P 必须 prefill 完整 ~50K tokens
+         ↓     即使 capacity-backup 开着，P 的 backup 只有 turn-0 的 ~1K
+         ↓     radix prefix cache 命中前 1K，剩余 49K 重算
+         ↓     实测耗时：~1.5-3s @ Qwen3-30B TP1
+         ↓
+t=~2000ms P 完成 prefill，KV 进入 mooncake transfer 队列
+            ↓
+t=~2050ms mooncake 开始 P→D' transfer
+         ↓
+         ↓ ←←← 大头 2：P→D mooncake transfer 段
+         ↓     KV 张量 ~5-9 GB（50K tokens × 2 bytes/token × layers × heads...）
+         ↓     **TCP loopback** 实测耗时：~1.5-4s
+         ↓     ↑↑↑ 当前 sweep 未启用 RDMA，走的是单机 lo 设备
+         ↓     若启用 IB RDMA @ 200 Gb/s，理论 200-400ms
+         ↓
+t=~4500ms transfer 完成，D' 上 session 重建好
+            ↓
+t=~4510ms D' 开始 decode（小幅度 append-prefill 余下的 ~2K append + 生成）
+            ↓
+t=~4550ms 首个 token 出来 → TTFT 测点
+```
+
+**单次 reseed 总耗时：3-7s**（中位 ~2.5s 来自较小 session，p99 ~7.7s 来自最大 session）。**re-prefill 段与 transfer 段大致五五开**，受 session 大小影响。
+
+### 这就是为什么 v2 的 TTFT p99 = 1.28s
+
+8.3% slow path 走的是上面这条流水线，其中 reseed 路径（`pd-router-d-session-reseed`）单独占 3.4%（150/4449 请求），构成 KVC TTFT p99 长尾的主要贡献。
+
+---
+
+## 3. 已审查的所有"看起来像 D→P 但其实不是"的代码
+
+下面这些在搜索时容易误判成 D→P 实现，**全部经独立 audit 排除**：
+
+| 文件:行 | 看起来像 | 实际是 |
+|---|---|---|
+| `replay.py:1483 _commit_prefill_backup_residency` | "把 backup 提交到 P" | bookkeeping 函数，更新 `session.prefill_resident_tokens` 计数字段。不传输任何 KV 数据，只在 seed/reseed 完成后被调用。 |
+| `replay.py:1572 _reserve_prefill_backup_capacity` | "预留 backup 空间" | 检查 P 池可用空间并按 LRU 驱逐别的 backup session 腾位置。不传 KV，只调整 reservation 计数。 |
+| `cli.py:182 --kvcache-prefill-backup-policy` | "backup 策略" | 只决定 reseed 完成后是否 `_close_prefill_session`。capacity-backup = 保留 P 端 streaming session 不关；release-after-transfer = 立刻关闭。**两种策略下 P 的 KV 都是 seed-time 的静态快照**。 |
+| `session_aware_cache.py:release_session` | "释放 session（可能含外发）" | 仅调 `kv_pool_allocator.free(kv_indices)`。零网络调用。 |
+| `disaggregation/decode.py: start_decode_thread` | "decode 端线程，可能有出站" | 纯 receiver loop。处理入站 `AUX_DATA / CHUNK_READY / STAGING_REQ / KVPoll.Success`，**没有出站 KV 传输分支**。 |
+| `disaggregation/mooncake/conn.py:1563` | "传输请求添加" | `assert disaggregation_mode == PREFILL`——硬约束，只有 P 端能调。 |
+| `mooncake.MooncakeKVSender` / `MooncakeKVReceiver` | "双向 sender / receiver" | 强角色化：Sender 只在 PREFILL 模式实例化，Receiver 只在 DECODE 模式。`BaseKVManager` 抽象无 bidirectional slot。 |
+| `pd-router-d-session-reseed-after-eviction` execution_mode | "走 backup 的快路径" | 实际还是走完整 `_invoke_kvcache_seeded_router`（P 完整 prefill + 完整 mooncake transfer），只是 `_eviction_suffix()` 在 execution_mode 字符串末尾加了 "-after-prefill-backed-eviction" 标签。**没有任何 fast-path 优化**。v2 中仅 2/4449 请求走到这个标签。 |
+
+---
+
+## 4. D→P 增量同步：要做的是什么
+
+完整 D→P 增量同步的设计目标：**让 P 端的 backup KV 在 direct-to-D append 完成后异步追上 D 端的 KV，让 reseed 退化为单次 P→D transfer（无需 P re-prefill）**。
+
+### 抽象数据流
+
+```
+当前：
+  direct-to-D append: D 本地 append-prefill，P 端 backup 锁住不变
+  reseed:             P re-prefill 完整 50K + P→D transfer 完整 50K
+
+目标：
+  direct-to-D append: D 本地 append-prefill，**同时**异步把新增的 KV 块推回 P
+  reseed:             P→D' transfer 完整 50K (already up-to-date)
+                      无需 P re-prefill
+```
+
+### 实现层面要改的事
+
+按工程难度排序：
+
+#### 4.1 Mooncake 双角色化（中等难度，~400 LOC）
+
+- `BaseKVSender` / `BaseKVReceiver` 抽象保留，但允许同一 worker 同时实例化两种角色
+- `MooncakeKVManager.__init__` 把 PREFILL / DECODE 分支改成"role set"，允许 worker 同时持有 sender 和 receiver
+- 新增 `DecodeKVSender` 类（D 端用于把 append KV 推回 P）
+- 新增 `PrefillKVReceiver` 类（P 端用于接收 D 的 append KV）
+- 引入第二个 bootstrap channel（避免与原 P→D 通道在 buffer pointer 协商上冲突）
+
+#### 4.2 D 端 append commit hook（容易）
+
+- 每次 `direct-to-D-session` 完成后，识别新写入的 KV 块（D scheduler 在 commit 时知道）
+- 入队 D→P 传输（异步，不阻塞 next request）
+- 标记 backup 是否成功送达 P（用于后续 reseed 决策）
+
+#### 4.3 P 端 radix tree 多生产者扩展（**最难，工程量主体**）
+
+**这是真正的架构 blocker**。SGLang 的 P 端 radix cache 当前假设：
+- 单一生产者（本 worker 的 model 输出）
+- 树插入只在 prefill / decode 完成时发生
+- KV 索引由本 worker 的 token_to_kv_pool_allocator 分配
+
+要让 P 接收 D 喂来的 KV 块，需要：
+- 扩展 radix tree 节点的写入路径，允许"外部供给的 KV + token 序列"被插入
+- 处理 KV 索引重映射（D 的 slot 号在 P 上无意义）
+- 处理 reference counting（同一 session 可能既被本 worker 用、又被 D 喂回更新）
+- 处理 eviction policy 协调（P 端 radix LRU 不应让"被 D 喂入的 backup"先被驱逐）
+- 处理 KV 数据格式的跨 worker 兼容（同样的 model layout，应该是 trivial，但需要测试）
+
+#### 4.4 agentic-pd-hybrid 端 hook（容易）
+
+- `_invoke_session_direct` 完成后，新增一步：触发 D→P 同步 RPC（异步）
+- `_invoke_kvcache_seeded_router` 在 reseed 触发前先 probe P 是否有 up-to-date backup；若有，跳过 re-prefill，只做 P→D transfer
+- 新增 CLI flag `--enable-d-to-p-sync`，默认 off，保留 baseline 行为
+- 新增 structural log channel 记录 D→P 同步事件 / 失败 / 延迟
+
+### 实现完毕后的预期收益
+
+| 指标 | 当前 (v2) | RDMA only | RDMA + D→P sync |
+|---|---:|---:|---:|
+| reseed re-prefill 段 | 1.5-3s | 1.5-3s（不变） | **~0**（已有 up-to-date backup） |
+| reseed transfer 段 | 1.5-4s | 0.2-0.4s | 0.2-0.4s |
+| reseed 总耗时 | 3-7s | 1.7-3.4s | **0.2-0.4s** |
+| TTFT p99 | 1.285s | ~0.7s | **~0.4-0.5s**（与 DP 接近或胜过） |
+| 8.4% slow path 占比 | 不变 | 不变 | 可能保持但单次代价大幅下降 |
+
+→ 这就是 paper 里 future-work 应当声明的**"完整版 KVC 才能真正在 TTFT 全分位数上击败 DP"** 的路径。
+
+---
+
+## 5. 仓库分支审查（确认无作者私下实现）
+
+`git ls-remote origin --refs` 完整结果：
+
+```
+9ccd853...  refs/heads/kvc-debug-journey-v1-to-v4   ← 本工作分支（含本文档）
+e9062b1...  refs/heads/main                          ← baseline，落后我们 18 commit
+```
+
+- **服务器只有 2 个分支**，**0 个 tag**，**0 个隐藏 ref**
+- main 是更老的 baseline；含 `_commit_prefill_backup_residency` 等同名函数，但语义与本工作分支一致——都是静态 backup，无 D→P 同步
+- 全 git 历史搜索 `D->P / d-to-p / decode.*prefill.*transfer / kv.*pushback / kv.*sync / incremental / mirror` 关键词，**唯一命中是 commit `9ccd853`**（本文档相关的 doc 改动）
+- 唯一 remote 是 `origin`（`git@ipads.se.sjtu.edu.cn:wangjh/agentic-pd-hybrid.git`），无 upstream / fork
+
+→ **作者没有在其它分支偷偷实现 D→P**。这块工作是真空。
+
+---
+
+## 6. 下一步
+
+按 ROI 排序：
+
+### 必做（落地下一阶段）
+
+1. **新开 `feat/d-to-p-sync` 分支** 从当前 `kvc-debug-journey-v1-to-v4` 起步
+2. 写设计文档 `docs/D_TO_P_SYNC_DESIGN_ZH.md`：
+   - 包括上面 §4 的实现细节
+   - 添加 sequence diagram（P/D 通信时序）
+   - 评估 SGLang radix tree 多生产者扩展的具体 API 改动
+   - 评估 D→P 同步对 direct-to-D fast path 自身延迟的影响（理想是异步零开销）
+3. POC 阶段 1：mooncake 双角色化 + 一个能跑通的 D→P transfer 单测
+4. POC 阶段 2：P 端 radix tree 多生产者扩展（重点工程量）
+5. POC 阶段 3：agentic-pd-hybrid 端的 hook + flag
+6. 端到端验证：跑同 trace 同 ts=1 配置，目标 TTFT p99 < 0.5s
+
+### 推荐
+
+7. **同时启用真 RDMA**（独立于 D→P 工作，只需改 sweep 脚本加 `--force-rdma --ib-device mlx5_0`），先把现有 transfer 段加速作为 baseline
+8. **跑 RDMA-only 对照**：先证明单 RDMA 启用能把 TTFT p99 从 1.28s 压到 ~0.7s，再用 D→P sync 把剩下的 re-prefill 段也吃掉。这样 paper 里能写两条独立的 ablation
+
+### 不要做的事
+
+- 在 main / 工作分支上做 D→P 实验（隔离开），主分支应该保持 v2 稳定
+- 试图通过 capacity-backup 现有 flag "调出"D→P 效果——它结构上做不到
+
+---
+
+## 附录 A：本文档涉及的代码位置
+
+| 函数 / 字段 | 位置 |
+|---|---|
+| `_commit_prefill_backup_residency` | `src/agentic_pd_hybrid/replay.py:1483` |
+| `_reserve_prefill_backup_capacity` | `src/agentic_pd_hybrid/replay.py:1572` |
+| `_close_prefill_session` | `src/agentic_pd_hybrid/replay.py:1507` |
+| `_close_decode_session` | `src/agentic_pd_hybrid/replay.py:1539` |
+| `_invoke_session_direct` (direct-to-D 路径) | `src/agentic_pd_hybrid/replay.py:2719` |
+| `_invoke_decode_session_direct` | `src/agentic_pd_hybrid/replay.py:2826` |
+| `_invoke_kvcache_seeded_router` (reseed 路径) | `src/agentic_pd_hybrid/replay.py:2101` |
+| `DirectSessionState.prefill_resident_tokens` | `src/agentic_pd_hybrid/replay.py:128` |
+| `_eviction_suffix` | `src/agentic_pd_hybrid/replay.py:1220` |
+| `--kvcache-prefill-backup-policy` CLI flag | `src/agentic_pd_hybrid/cli.py:182-189, 436-441` |
+| `MooncakeKVManager.__init__` | `third_party/sglang/python/sglang/srt/disaggregation/mooncake/conn.py:187-256` |
+| `start_decode_thread` (decode 端 receive loop) | `third_party/sglang/python/sglang/srt/disaggregation/mooncake/conn.py:1425-1496` |
+| `add_transfer_request` (assert PREFILL) | `third_party/sglang/python/sglang/srt/disaggregation/mooncake/conn.py:1563` |
+| `MooncakeKVSender` / `MooncakeKVReceiver` | `third_party/sglang/python/sglang/srt/disaggregation/mooncake/conn.py:1648, 1740` |
+| `BaseKVSender` / `BaseKVReceiver` 抽象 | `third_party/sglang/python/sglang/srt/disaggregation/base/conn.py` |
+| `session_aware_cache.release_session` | `third_party/sglang/python/sglang/srt/mem_cache/session_aware_cache.py:250-276` |
+| `session_controller._close` | `third_party/sglang/python/sglang/srt/managers/session_controller.py:293-316` |
+
+## 附录 B：相关 commit
+
+| Commit | 内容 |
+|---|---|
+| `9ccd853` | docs: D→P 缺口的 Opus forensic audit 写入 V2_DEEP_ANALYSIS §4.2 + KVC_ROUTER_ALGORITHM §9 |
+| `2ec0deb` | v2 实现（reset-on-success + threshold 2048→8192）—— 直接 trigger 了对 reseed 慢路径的关注 |
+| `c47adaf` | feat: backpressure pause hint（与 reseed 不直接相关，但展示了"D 端可主动告知 router"的通信通道存在，是未来 D→P sync 控制平面的潜在基础） |
+
+## 附录 C：相关 paper 章节建议
+
+- **§Background**：把 §1-§2 的 reseed 现状作为 motivation 摆出
+- **§Algorithm**：参考 `KVC_ROUTER_ALGORITHM.md` Algorithm 1-3
+- **§Evaluation §Slow Path Cost**：把 §2 的端到端时间线作为 Figure（sequence diagram）
+- **§Future Work / Limitations**：把本文 §4 作为 KVC 真正实现"完整 fast path 替代"的 roadmap，引用 D→P 工作的设计文档（后续 `feat/d-to-p-sync` 分支产物）
+
+---
+
+**核心句**：v2 实现的 KVC 在 91.6% 请求上证明了 session-affinity 路由的价值，但 8.3% 的 reseed 慢路径让 TTFT p99 比 DP 差 3×。这条慢路径的 50% 时间在 P 端 re-prefill、50% 在 mooncake transfer——RDMA 只能救后者，**D→P 增量 KV 同步是唯一能消除 re-prefill 的机制**，且当前在框架、SGLang、mooncake 三层都没有实现，需要新建 `feat/d-to-p-sync` 分支从设计文档开始。

docs/STRUCTURAL_VALIDATION_REPORT_ZH.md

@@ -0,0 +1,304 @@
+# 结构性缺陷验证报告
+
+**日期**：2026-05-06
+**对照数据源**：
+- `outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/`（v5 KVC kv-aware Option D，2P6D，**3 次同配置 rerun**）
+- `outputs/qwen3-30b-tp1-exps/exp1_8way_dp_cache_aware_summary.json`（同 trace 8DP CA）
+- `outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/.../logs/decode-{0..5}.log`、`prefill-{0,1}.log`
+**模型**：Qwen3-30B-A3B（TP1），单机 8×H100 80GB，trace `qwen35-swebench-50sess.jsonl`（4449 reqs / 52 sessions）。
+**报告作用域**：验证 `docs/AGENTIC_FIT_ANALYSIS_ZH.md` §1-§7 提出的结构性 claim 是否真实存在；量化影响。
+
+> ⚠️ **环境限制**：本轮缺 GPU 访问，未跑新 sweep。所有数据来自已存在的 v5 rerun + 8DP baseline。Backpressure 代码已实现但**未端到端验证**——下文标注为"预期收益（pending GPU smoke）"。
+
+---
+
+## 0. 实验有效性锚点：N=1 不可信
+
+3 次 v5 baseline EXP2（**完全相同配置**）的 errors 漂移：
+
+| Run | Errors | Lat P50 | Lat P90 | TTFT P50 |
+|---|---:|---:|---:|---:|
+| run1 | **372** | 1.11s | 8.65s | 0.147s |
+| run2 | **912** | 0.94s | 7.68s | 0.071s |
+| run3 | **396** | 1.22s | 8.43s | 0.183s |
+
+errors 漂移 **2.5×**（372 → 912），P50 latency 漂移 **30%**。**任何 N=1 比较 < 30% 差异都不可信。** 后续所有"同 trace 不同配置 / 不同代码"的对比，都需要 N≥3 才有意义。
+
+**对 KVC vs DP 的 headline 数据，3 次 KVC 的最佳值（P50=0.94s）仍然是 DP（P50=0.65s）的 1.45×**——8 way DP 的优势远超 single-run variance 范围，这一头条结论不受 variance 影响。
+
+---
+
+## §1. Session 永久 pin 到 D + 容量盲选 → 极端双峰 ✅ 完全成立
+
+### Claim
+KvAwarePolicy 评分以 hash overlap 为主，没有 D 容量项。Session 第一次落到某 D 后被永久 pin。导致大 session 在已满 D 上反复 admission 拒绝，小 session 在原 D 上 100% 走 direct-to-D。
+
+### 数据
+
+**(a) Session 永久绑定，跨 3 次 rerun 一致**：
+
+```
+run1: 52 sessions, avg distinct-D-per-session = 1.00
+run2: 52 sessions, avg distinct-D-per-session = 1.00
+run3: 52 sessions, avg distinct-D-per-session = 1.00
+```
+
+每个 session 在整个运行中只访问 **1 个** D worker，3 次独立 run 完全一致。**不是巧合，是结构。**
+
+**(b) Direct-to-D 命中率呈极端双峰**：
+
+| Direct-to-D rate | run1 | run2 | run3 |
+|---|---:|---:|---:|
+| 0-20%（饿死） | 15 | 18 | 16 |
+| 20-40% | 7 | 6 | 7 |
+| 40-60% | 11 | 7 | 9 |
+| 60-80% | 5 | 6 | 4 |
+| 80-100%（顺利） | 14 | 15 | 16 |
+
+中间态稀少，两端拥挤。
+
+**(c) 跨 3 次 run 一致饿死的 session 与 session 大小强相关**：
+
+```
+13 sessions starved (<20% direct-to-D) in ALL 3 runs.
+  avg peak input of consistently-starved sessions: 62043 tokens
+  avg peak input of consistently-lucky sessions:   31344 tokens
+  ratio: 1.98× — starved sessions are exactly 2× larger.
+```
+
+**13/52 = 25% 的 session 在 3 次独立 run 中都被饿死，且这些 session 的 peak input 恰好是顺利 session 的 2 倍。** 这排除了"运气"假说，证实是大 session 在容量过载 D 上结构性失败。
+
+### 影响量化
+- 25% session 几乎每个 turn 都走 fallback 路径，相对 direct-to-D **TTFT 慢 100×、E2E 慢 6×**（数据点：fallback path mean lat ~3.5s vs direct ~0.5s）
+- 对应这些 session 的用户体验是"系统性糟糕"，而不是"偶尔慢"
+- **SLO 视角下 P99 完全由这 13 个 session 拉高**
+
+### 结论
+**完全成立**。修复方向（不在本轮）：policy score 加 capacity penalty + 允许 session 跨 D 迁移，或 D 端引入 hot session retract。
+
+---
+
+## §2. D 端 LRU 只 evict idle session → 跟不上压力 ✅ 完全成立
+
+### Claim
+`scheduler.py:2040` 的 `evict_idle_streaming_sessions_lru` 只能 evict "所有 req 都 finished + streaming 模式"的 session。高并发下 hot session 永远不 idle，LRU 找不到东西可踢。结果 D 顶到 100% 然后撞 mooncake transfer timeout。
+
+### 数据（v5 baseline rerun run1）
+
+| D worker | Trim 事件 | KVTransferError | 峰值 token_usage |
+|---|---:|---:|---:|
+| decode-0 | 9 | 0 | 0.99 |
+| decode-1 | 43 | 4 | 0.99 |
+| decode-2 | 16 | 153 | 0.97 |
+| decode-3 | 37 | 29 | 0.99 |
+| decode-4 | 28 | 90 | **1.00** |
+| decode-5 | 30 | 93 | **1.00** |
+
+**6 个 D 全部峰值 ≥ 0.97**，其中 2 个直接顶到 1.00（KV 池完全耗尽）。**LRU 触发 9-43 次，远不及 transfer 错误的 90-153 次。**
+
+decode-2 极端：trim 16 次 vs error 153 次 = LRU 比错误慢 **9.5×**。
+
+### 影响量化
+- 单 run 累计 369 KVTransferError（总 6 个 D 之和）
+- 对应 ~8% 的请求失败率（v5 errors 9/372/912 三次平均 ~430/4449 = 9.7%）
+- **每次 mooncake timeout 是 32s**——对 P99 latency 直接贡献几十秒尾巴
+
+### 结论
+**完全成立**。修复方向（不在本轮）：分层 eviction——除 idle 外加冷 session retract、按访问频率/时序加权。Backpressure（本轮代码）只是把"D 满"的雪崩从"timeout 错误"转成"主动等待"，**不是真正解决容量问题**。
+
+---
+
+## §3. 没有 D→Replay backpressure 通道 ✅ 成立（已实现修复）
+
+### Claim
+D 端 transfer queue 堆 → 32s timeout → KVTransferError，没有"D 过载请慢点"信号反向到 replay；concurrency 一直 32 不降。
+
+### 数据
+- §2 的 369 KVTransferError 全部为 32s mooncake timeout（日志中均为 `Failed to send kv chunk` 或 `Decode instance could be dead`）
+- 错误集中在运行后半段（按现有 `KVC_DEBUG_JOURNEY_V1_TO_V5.md` §v4：错误均在 run 的 44.8% 之后开始累积）
+- 表明：**前期 D 容量充裕时正常，达到容量上限后所有后续请求集中失败**——典型无 backpressure 系统行为
+
+### 修复（本轮已实现，待 GPU smoke 验证）
+
+代码改动：
+1. `third_party/sglang/python/sglang/srt/managers/io_struct.py`：`DirectAppendAdmissionReqOutput` 增加 `recommended_pause_ms` 字段
+2. `third_party/sglang/python/sglang/srt/managers/scheduler.py:admit_direct_append`：基于 `transfer_queue_depth`、`retracted_queue_depth`、`token_usage_after` 计算 hint
+   ```python
+   def _compute_backpressure_pause_hint(...):
+       if retracted_queue_depth > 0: return 1500
+       if token_usage_after >= 0.90: return max(200, min(2000, overshoot * 5))
+       if transfer_queue_depth >= 8: return min(2000, transfer_queue_depth * 100)
+       return 0
+   ```
+3. `src/agentic_pd_hybrid/replay.py`：
+   - `DecodeResidencyState.pause_until_s: dict[str, float]`
+   - `_query_decode_direct_admission` 解析 hint 更新 `pause_until_s`
+   - 新增 `_wait_for_decode_pause`，在 `_invoke_router` / `_invoke_session_direct` 入口检查
+4. CLI flag：`--enable-backpressure`、`--backpressure-max-pause-s 2.0`（默认关闭）
+5. 结构性日志：`structural/admission-events.jsonl`、`backpressure-events.jsonl`、`session-d-binding.jsonl`
+
+### 预期收益（pending GPU smoke E2 vs E1）
+- KVTransferError 应从 ~370 / 4449 跌到 < 50 / 4449
+- P99 应改善（消除 32s timeout 尾巴）
+- 整体 latency mean 可能**略升**（被强制 pause），但 P99 应大幅降
+- backpressure-events.jsonl 应显示 D-4 / D-5 累积大量 pause 事件（与 §2 数据吻合）
+
+### 结论
+**Claim 成立；修复已实现，待 smoke 验证**。注意：backpressure 是**降级**机制，不是性能优化——它把"硬错误"换成"主动等待"，整体 throughput 不会因此提升。
+
+---
+
+## §4. Admission RPC 与 scheduler 主循环耦合 ⚠️ 间接证据，本轮未直接验证
+
+### Claim
+`admit_direct_append` 进 scheduler 主循环遍历 session slot，admission RPC 频率 16+/s 时与 decode 抢调度。
+
+### 现有间接证据
+- `docs/V5_PROFILE_INVESTIGATION_ZH.md`：仅加 1Hz `/server_info` polling 就让 EXP2 errors 从 9 涨到 415（46×）；但 v6 P0 三次 baseline 不开 polling 同样得到 372/912/396——**polling 不是唯一原因，主循环负载本身就敏感**。
+
+### 本轮未做
+- 没有"admission probe 拆 fast/slow"的对照实验。需要 SGLang 较深的改动（提供 lock-free snapshot），不在 KISS 边界。
+
+### 结论
+**Claim 间接成立，本轮未直接验证**。Backpressure 实现里 admission RPC 的频率没有变（仍每个 turn 一次），只是结果会触发 sleep。如果这条 claim 成立，加 backpressure 后 admission RPC 数量大致不变但每次响应里的 `pause_ms` 会非零——**新增的 admission-events.jsonl 可在 GPU smoke 后用来直接验证此现象**。
+
+---
+
+## §5. P-side round-robin 不感知 D 健康 ✅ 成立
+
+### Claim
+`pd_router.py:_select_decode_index` 是裸 round-robin。任一 P 撞到 hot D 时反复失败，另一 P 完全不受影响。
+
+### 数据（v5 baseline rerun run1）
+
+| Worker | KVTransferError | "Decode could be dead" |
+|---|---:|---:|
+| prefill-0 | **367** | 361 |
+| prefill-1 | **2** | 0 |
+
+prefill-0 的请求量从 summary 看是 2225 vs prefill-1 的 2224——**请求量近乎对半，错误率差 180×**。
+
+### 影响量化
+- 失败请求集中在 P-0 → 某个 hot D 的链路上（日志中反复出现 `to 10.45.80.47:XXXXX`）
+- 单 P 的"死亡链路"贡献了 **99%** 的全部 KVTransferError
+- 如果 P 选择能避开"正在和 hot D 死磕"的链路，**理论上可消除单 P 故障的雪崩效应**
+
+### 备注
+- 此现象**未在 v6 P0 的 3 次 rerun 中横向验证**——只有 run1 的日志可读。需要在新 sweep 的 prefill-{0,1}.log 上重复确认，避免 N=1 嫌疑。
+
+### 结论
+**单 run 数据成立，多 run 一致性未验证**。修复方向（不在本轮）：router 选 P 时考虑 (P 当前 inflight transfer 数, 目标 D 健康度)。
+
+---
+
+## §6. （已撤回）Replay 端 session footprint 估算膨胀
+
+写计划时仔细看代码后撤回——`_estimate_session_resident_tokens` 返回 full prompt，但所有需要"增量"的 call site (`replay.py:1247-1254`、`:1393-1394`、`:1490-1491`) 都已用 `target - current` 减法处理。**不是 bug**。
+
+---
+
+## §7. time-scale=10 把 inter-turn gap 压到 1/10 ✅ 完全成立
+
+### 数据
+
+```
+原始 trace inter-turn gap (n=4397):
+  p10=1.6s   p50=2.5s   p90=7.8s   p99=25.1s   max=261s
+
+time-scale=10 实际 replay gap:
+  p10=0.16s  p50=0.25s  p90=0.78s  p99=2.5s    max=26s
+```
+
+真实 agentic 用户/agent 在 turn 之间停 2-8 秒（思考、打字、tool call、agent reasoning）。time-scale=10 把这些窗口压到 0.16-0.78 秒——**人为消除了 D 的自然 idle 时间**，正好是 KVC 想利用的"session 短暂 idle 时 LRU 可以 evict、其他 session 可以 admit"机会。
+
+### 测量学影响
+- 所有 v3-v6 数据基于 time-scale=10
+- 意味着所有"KVC 在 SWE 上输给 baseline"的结论**可能被 benchmark 放大了**
+- §1 的 25% session 永久饿死现象，在 time-scale=1 下可能因为 D 有更多 drain 时间而显著缓解
+
+### 本轮未做
+- 没跑 time-scale=1 baseline。这是项目当前**最重要但缺失的验证**。
+- Smoke sweep 脚本（`scripts/sweep_backpressure_smoke.sh`）E3、E4 包含了 time-scale=1 的 KVC + DP 短 trace 对比，等 GPU 时跑。
+
+### 结论
+**Claim 完全成立；time-scale=1 验证为 P0 待办**。
+
+---
+
+## 头条对比（同 trace、同硬件）
+
+```
+8-way DP cache-aware (TP1):
+  errors=  0 | latency mean=1.426s p50=0.654s p90=3.609s
+              | TTFT  mean=0.123s p50=0.093s p90=0.256s
+
+KVC v5 2P6D (3 reruns, no polling):
+  run1: errors=372 | mean=3.50s p50=1.11s p90=8.65s | TTFT mean=2.13s
+  run2: errors=912 | mean=3.00s p50=0.94s p90=7.68s | TTFT mean=1.64s
+  run3: errors=396 | mean=3.42s p50=1.22s p90=8.43s | TTFT mean=2.07s
+```
+
+KVC 三次 run 全输 DP，且差距远超 single-run variance：
+- Latency mean：DP 优 **+110%**（KVC 平均 3.30s vs DP 1.43s）
+- Latency P50：DP 优 **+65%**（KVC 平均 1.09s vs DP 0.65s）
+- TTFT mean：DP 优 **+1500%**（KVC 平均 1.95s vs DP 0.12s——慢 17×！）
+- Errors：DP 0 vs KVC 平均 ~560
+
+**这是这个项目当前最严肃的事实**——所有 KVC 复杂度回报为负。
+
+---
+
+## 综合结论
+
+按"是否结构性 + 影响大小"的二维分类：
+
+| Claim | 结构性 | 影响 | 本轮验证 | 修复（KISS 内） | 修复（KISS 外） |
+|---|---|---|---|---|---|
+| §1 Session pin + 容量盲选 | 强 | 大（25% session 饿死） | ✅ 3 run 一致 | ❌ | capacity-aware policy + 跨 D 迁移 |
+| §2 LRU 跟不上 | 强 | 大（每次 ~370 KVTransferError） | ✅ 6 D 数据 | ❌ | 分层 eviction、hot retract |
+| §3 无 backpressure | 强 | 中-大（消除 32s timeout 雪崩） | ⚠️ 已实现，待 smoke | ✅ **本轮交付** | – |
+| §4 admission RPC 干扰 | 弱-中 | 中 | ⚠️ 间接 | ❌ | probe / commit_evict 拆分 |
+| §5 P-side 不感知 D 健康 | 中 | 中（单 P 错误率差 180×） | ✅ N=1，需 N≥3 复核 | ❌ | router P 选择带 D 健康反馈 |
+| §6 estimate 膨胀 | – | – | ❌ 已撤回 | – | – |
+| §7 time-scale=10 失真 | 强（测量学） | 大（可能颠覆所有 KVC vs DP 结论） | ✅ 数据明确 | ✅ 改 flag | – |
+
+### 最关键的两个 takeaway
+
+1. **§7 time-scale=1 是当前项目所有结论的前置依赖**——必须先做。如果 time-scale=1 下 KVC 与 DP 接近，前面所有 v3-v6 的"KVC 输得彻底"诊断都需要重新解读。
+2. **§1 + §2 是双胞胎结构性问题**——session 被永久 pin 在某个 D + D 不能 evict 已满 = 大 session 永久卡死。任何不动 policy + 不动 LRU 的修复（包括本轮的 backpressure）只能让症状好看，不能消除根因。
+
+---
+
+## 本轮代码改动汇总（git diff 范围）
+
+```
+src/agentic_pd_hybrid/replay.py        # +结构性日志 + backpressure pause 检查 + admission 增强
+src/agentic_pd_hybrid/cli.py           # +CLI flags
+src/agentic_pd_hybrid/benchmark.py     # +CLI flags 透传
+third_party/sglang/python/sglang/srt/managers/io_struct.py
+third_party/sglang/python/sglang/srt/managers/scheduler.py
+                                       # +recommended_pause_ms 字段 + hint 计算
+scripts/sweep_backpressure_smoke.sh    # 4-run smoke sweep（待 GPU 跑）
+scripts/analysis/analyze_backpressure_smoke.py
+                                       # 配套分析器
+docs/REFACTOR_PLAN_ZH.md               # 计划文档
+docs/STRUCTURAL_VALIDATION_REPORT_ZH.md
+                                       # 本报告
+```
+
+代码默认行为**不变**（`enable_backpressure=False`）——所有现有脚本/配置无影响。
+
+---
+
+## 待 GPU 时执行
+
+```bash
+bash scripts/sweep_backpressure_smoke.sh
+python3 scripts/analysis/analyze_backpressure_smoke.py outputs/sweep_backpressure_smoke
+```
+
+预算：4 个 run × 30-60 min ≈ 3-4h GPU 时间。
+
+按 §3 的预期：E2 (KVC + backpressure) 相对 E1 (KVC baseline) 应有 errors 降 70%+；P99 改善；TTFT P50 持平或略升。E3 (KVC + backpressure @ time-scale=1) vs E4 (DP @ time-scale=1) 是验证 §7 的关键对照。
+
+如果 E2 vs E1 的 errors 没有显著下降，说明 backpressure hint 公式调得不对（`_compute_backpressure_pause_hint` 阈值可调），或 §3 实际不是雪崩主因（更可能是 §2 D-side LRU 才是）。

docs/SWEBENCH_EXPERIMENT_PROGRESS.md

@@ -0,0 +1,95 @@
+# SWE-Bench PD Hybrid Experiment Progress
+
+## 实验目标
+
+在单节点 8xH100 上复现 agentic-pd-hybrid 三种 serving mechanism，对比 Qwen3.5-35B-A3B 在 SWE-Bench 500 instance agentic trajectory 上的性能。
+
+## 硬件环境
+
+- 8x H100 80GB (NVLink 互联, 2 NUMA nodes: GPU 0-3 / GPU 4-7)
+- 无 RDMA/IB 设备
+- Transfer backend: **mooncake TCP** (nixl UCX 因 pip 包缺少 CUDA 支持导致 segfault，已放弃)
+
+## 实验矩阵
+
+| 实验 | Mechanism | Workers | GPU 分配 | Router | Policy |
+|------|-----------|---------|----------|--------|--------|
+| A | pd-disaggregation | 1P + 1D (TP4 each) | P: 0-3, D: 4-7 | Yes | default |
+| B | pd-colo | 2 direct (TP4 each) | D0: 0-3, D1: 4-7 | No | default |
+| C | kvcache-centric | 1P + 1D (TP4 each) | P: 0-3, D: 4-7 | Yes | default |
+
+## 测试负载
+
+- 源数据: `simm-swe-bench/outputs/20260416-205833-hicache-qwen35-verified-0-500/audit.jsonl`
+- 39,417 lines (turns), 497 unique instances (sessions)
+- 每个 instance 8-150 turns (均值 79.3)
+- 转换为 agentic-pd-hybrid trace 格式: `outputs/qwen35-swebench-500.jsonl`
+
+## 关键发现
+
+### Transfer Backend 选择
+
+- **nixl (UCX)**: pip 安装的 nixl_cu12 包自带的 UCX 库没有 CUDA 支持，导致 GPU memory registration 时 segfault。系统 UCX (/opt/hpcx/ucx) 有 CUDA 支持但因 RPATH 无法被 NIXL 使用。
+- **mooncake (TCP)**: 可用。需要两处修改:
+  1. `third_party/sglang/.../mooncake_transfer_engine.py`: 从环境变量 `MOONCAKE_PROTOCOL` 读取协议，而非硬编码 `"rdma"`
+  2. `src/agentic_pd_hybrid/stack.py`: 当 `transfer_backend == "mooncake"` 且非 `force_rdma` 时，自动设置 `MOONCAKE_PROTOCOL=tcp`
+
+### 代码修改记录
+
+1. **`third_party/sglang/python/sglang/srt/distributed/device_communicators/mooncake_transfer_engine.py`**
+   - 将 `"rdma"` 硬编码改为 `os.environ.get("MOONCAKE_PROTOCOL", "rdma")`
+
+2. **`src/agentic_pd_hybrid/stack.py`**
+   - 在 `_build_process_env()` 中添加: mooncake 非 force_rdma 时默认设置 `MOONCAKE_PROTOCOL=tcp`
+
+3. **`scripts/convert_audit_to_trace.py`** (新建)
+   - 将 sibench audit.jsonl 转换为 agentic-pd-hybrid trace 格式
+
+## 实验进度
+
+- [x] Step 0: 环境准备 (uv sync, nixl/mooncake 安装)
+- [x] Step 1: Trace 格式转换 (39,417 lines 验证通过)
+- [x] Step 2: Smoke test (pd-disaggregation, mooncake TCP, 100 requests) — **通过**
+  - 100/100 requests, 0 errors
+  - Mean latency: 1.53s, P50: 0.77s, P90: 2.82s
+  - TTFT: mean 0.49s, P50 0.29s; TPOT: mean 4.7ms
+  - 91/100 cache hits
+- [x] Step 3a: 实验 A 全量尝试 (39K reqs, 497 sessions) — **中止**
+  - Run dir: `outputs/swebench-exps/pd-disaggregation-default-20260426T171113Z` (无metrics,被kill)
+  - 前 90% 完成 ~80min (~8-10 req/s), 但尾部 D 侧 KV cache 98% 饱和
+  - 497 并发 session 争抢 D 侧 token 空间, mamba 80-93 sessions 无法 drain
+  - **教训**: 1P+1D (TP4) 无法支撑 497 并发 session, 需减少 session 数量或降低 concurrency
+- [x] Step 3b: 实验 A — pd-disaggregation (52 sessions, 4449 reqs, concurrency=32) — **完成**
+  - Run dir: `outputs/swebench-exps/pd-disaggregation-default-20260426T202540Z`
+  - Trace: `outputs/qwen35-swebench-50sess.jsonl` (10% sample, 52 sessions)
+  - **结果**: 4449/4449 成功, 0 errors
+  - Latency: mean=1.66s, P50=0.97s, P90=3.64s, P99=7.68s
+  - TTFT: mean=0.45s, P50=0.34s, P90=0.88s
+  - TPOT: mean=5.2ms, P50=5.2ms
+  - Cache hit: 4199/4449 (94.4%)
+- [x] Step 4: 实验 B — pd-colo — **失败: SGLang bug**
+  - Run dir: `outputs/swebench-exps/pd-colo-default-20260426T210129Z`
+  - **Bug**: `--disaggregation-mode null` (colocation) 下 Qwen3.5-35B-A3B 模型触发 token_to_kv_pool_allocator 内存泄漏
+  - 错误: `ValueError: token_to_kv_pool_allocator memory leak detected!`
+  - 两个 direct worker 在处理 ~5 个请求后均 crash (Scheduler exception)
+  - **结论**: 当前 vendored SGLang v0.5.10 不支持 Qwen3.5-35B-A3B 的 colocation 模式
+- [x] Step 5: 实验 C — kvcache-centric — **完成 (高错误率)**
+  - Run dir: `outputs/swebench-exps/kvcache-centric-default-worker-admission-20260426T210800Z`
+  - 4390/4449 errors (98.7%) — admission control 过于保守
+  - 59 成功请求: mean latency 1.24s (比 pd-disagg 快 25%), TTFT 0.18s (快 60%)
+  - 详细分析见 `docs/SWEBENCH_EXPERIMENT_RESULTS.md`
+- [x] Step 6: 结果对比分析 — **完成**
+  - 完整报告: `docs/SWEBENCH_EXPERIMENT_RESULTS.md`
+
+## 启动脚本
+
+- `scripts/run_exp_a_pd_disagg.sh` — 实验 A
+- `scripts/run_exp_b_pd_colo.sh` — 实验 B
+- `scripts/run_exp_c_kvcache_centric.sh` — 实验 C
+- `scripts/convert_audit_to_trace.py` — Trace 转换
+
+## 已知风险
+
+1. Qwen3.5-35B-A3B TP4 可用 mem ~12GB/GPU (after model + CUDA graph)，长 session (150 turns) 可能 OOM
+2. mooncake TCP loopback 延迟远低于真实跨机，结果偏乐观
+3. 原始 trace 时间跨度 ~6000s，全量回放非常耗时

docs/SWEBENCH_EXPERIMENT_RESULTS.md

@@ -0,0 +1,121 @@
+# SWE-Bench PD Hybrid Experiment Results
+
+## 实验配置
+
+- **模型**: Qwen3.5-35B-A3B (MoE, 35B total / 3B active), TP4
+- **硬件**: 8x H100 80GB, NVLink, 单节点
+- **Transfer backend**: mooncake TCP (loopback)
+- **Trace**: 52 sessions, 4,449 requests (10% sample of SWE-Bench 500 instances)
+- **时间压缩**: time-scale=10, concurrency-limit=32
+
+## 结果汇总
+
+### Experiment A: pd-disaggregation (baseline)
+
+| Metric | Value |
+|--------|-------|
+| Run dir | `pd-disaggregation-default-20260426T202540Z` |
+| Requests | 4,449 / 4,449 (100%) |
+| Errors | 0 |
+| **Mean Latency** | **1.662s** |
+| P50 Latency | 0.973s |
+| P90 Latency | 3.644s |
+| P99 Latency | 7.676s |
+| Mean TTFT | 0.445s |
+| P50 TTFT | 0.340s |
+| P90 TTFT | 0.880s |
+| Mean TPOT | 5.20ms |
+| Cache Hit Rate | 94.4% (4199/4449) |
+| Mean Cached Tokens | 27,794 |
+| KV Transfer Blocks | 105,235 |
+
+### Experiment B: pd-colo (colocation) — FAILED
+
+| Metric | Value |
+|--------|-------|
+| Run dir | `pd-colo-default-20260426T210129Z` |
+| Status | **CRASHED** |
+| Error | `token_to_kv_pool_allocator memory leak detected!` |
+| Root Cause | SGLang v0.5.10 `--disaggregation-mode null` 与 Qwen3.5-35B-A3B (Mamba/GDN hybrid) 不兼容 |
+| Requests | ~10 / 4,449 (0.2%) |
+
+**结论**: 当前 vendored SGLang 不支持此模型的 colocation 模式。需要修复 token_to_kv_pool_allocator 中 Mamba 模型的内存管理。
+
+### Experiment C: kvcache-centric (session-aware PD)
+
+| Metric | Value |
+|--------|-------|
+| Run dir | `kvcache-centric-default-worker-admission-20260426T210800Z` |
+| Requests | 4,449 total |
+| **Errors** | **4,390 (98.7%)** |
+| Successful | 59 (1.3%) |
+| Mean Latency (success) | 1.238s |
+| P50 Latency (success) | 0.484s |
+| P90 Latency (success) | 2.550s |
+| Mean TTFT (success) | 0.179s |
+| P50 TTFT (success) | 0.081s |
+| Mean TPOT (success) | 4.70ms |
+| Direct-to-D Sessions | 56 |
+| KV Transfer (actual) | 196 blocks (vs 105,235 planned) |
+
+**Execution Mode 分布**:
+- `kvcache-centric` (failed): 4,390
+- `kvcache-direct-to-d-session` (success): 56
+- `pd-router-*` variants: 3
+
+## 关键分析
+
+### 1. pd-disaggregation (A) — 稳定可靠
+
+- 100% 成功率，0 错误
+- Mean latency 1.66s 合理 (包含 P→D KV transfer 开销)
+- 94.4% cache hit 说明 prefix cache 在 P 侧工作良好
+- KV transfer 105K blocks = 主要开销来源
+- **适合生产使用**
+
+### 2. pd-colo (B) — 不可用
+
+- Qwen3.5-35B-A3B 的 Mamba/GDN hybrid 架构在 `disaggregation-mode null` 下触发内存泄漏
+- 这是 SGLang 的 bug，不是 agentic-pd-hybrid 的问题
+- **需要 SGLang 修复后重新测试**
+
+### 3. kvcache-centric (C) — Admission 过于保守
+
+- 98.7% 错误率说明 admission control 拒绝了几乎所有请求
+- `kvcache-seed-min-turn-id=2` 过滤了 turn 1 的 seed（正确行为）
+- 但绝大多数 turn 2+ 请求也走 `kvcache-centric` 模式后失败
+- 可能原因:
+  - Worker admission 查询发现 D 侧没有对应 session 的 KV cache（因为 turn 1 没有 seed）
+  - D 侧 transfer queue 积压导致 admission 拒绝
+- 成功的 56 个 `direct-to-d-session` 请求表现优异: TTFT 0.08s (P50), 比 pd-disagg 的 0.34s 快 4x
+- **需要调优 admission 参数，或使用 `kvcache-seed-min-turn-id=1` 允许 turn 1 seed**
+
+### 4. kvcache-centric 成功请求 vs pd-disaggregation 对比
+
+| Metric | pd-disagg (A) | kvcache-centric (C, success only) | Delta |
+|--------|:---:|:---:|:---:|
+| Mean Latency | 1.662s | 1.238s | **-25.5%** |
+| P50 Latency | 0.973s | 0.484s | **-50.3%** |
+| Mean TTFT | 0.445s | 0.179s | **-59.8%** |
+| P50 TTFT | 0.340s | 0.081s | **-76.2%** |
+| Mean TPOT | 5.20ms | 4.70ms | -9.6% |
+| Actual KV Transfer | 105,235 blk | 196 blk | **-99.8%** |
+
+**当 kvcache-centric 成功时，性能提升显著：**
+- TTFT 降低 60-76% (D 侧直接 append，无需 P→D transfer)
+- 端到端 latency 降低 25-50%
+- KV transfer 减少 99.8%
+
+## 后续建议
+
+1. **修复 pd-colo**: 提交 SGLang issue 关于 Mamba/GDN 模型在 disaggregation-mode null 下的内存泄漏
+2. **调优 kvcache-centric admission**:
+   - 尝试 `--kvcache-seed-min-turn-id 1` 允许 turn 1 seed
+   - 放宽 `--kvcache-seed-max-decode-transfer-queue-reqs` 阈值
+   - 使用 `--kvcache-admission-mode router` (shadow state, 不在 critical path)
+3. **增加 D 侧内存**: 调整 `--mem-fraction-static` 给 KV cache 更多空间
+4. **多 P/D 配置**: 测试 2P2D (TP2) 配置以增加并行度
+
+## 实验日期
+
+2026-04-27

docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md

@@ -0,0 +1,641 @@
+# agentic-pd-hybrid 现框架性能与结构性问题报告
+
+**对象**：项目团队同学
+**前置假设**：读者**没看过** v3-v6 KVC 实验日志
+**数据范围**：项目仓库 `outputs/` 下截止 2026-05-06 的全部实验产物
+**目的**：把"现状"和"问题"分别交代清楚，给后续改造提供共同事实基础
+
+---
+
+## 0. 给没看过实验的读者：基础概念速览
+
+### 0.1 项目目标
+验证 **session-aware / KV-cache-aware P/D routing** 在 **agentic coding workload**（多轮 session、长 context、增量 append）上能否降低端到端延迟。基线对比对象是 vanilla SGLang xPyD。
+
+### 0.2 三种部署机制（**这三个名词全程会用**）
+
+| 机制 | 形态 | KV 流向 |
+|---|---|---|
+| **pd-disaggregation**（"PD disagg"） | P 和 D 是独立进程、分占不同 GPU | 每个请求 P 算 prefill → mooncake 推 KV → D 解码 |
+| **pd-colo**（"DP"，data-parallel） | 没有 PD 拆分，N 个独立完整 worker（每个自己 prefill+decode） | 没有 KV transfer；router 按 hash 分配请求 |
+| **kvcache-centric**（"KVC"） | 部署形态同 PD disagg；**D 上多了 SessionAwareCache**，能跨 turn 保留 session KV | 运行时决策：可走 direct-to-D（无 P）、可走 P→D disagg、可走带 reseed 的混合 |
+
+**Direct-to-D**（"D-direct"）：KVC 的快路径——D 上已有该 session 的 KV，新 turn 在 D 本地做 append-prefill，零 P 介入、零 mooncake transfer。这是 KVC 理论上能省时间的核心。
+
+**Fallback**：KVC admission 拒了 / 阈值不满足 / D 不健康时，退化到普通 PD disagg 路径。
+
+**Routing policy**（与机制正交）：
+- `default`：纯 round-robin
+- `sticky`：turn 2+ 黏到 session 的 last D
+- `kv-aware`：按 hash overlap + sticky 评分选 D（**KVC 必须配它**才能正确工作）
+
+### 0.3 数据来源
+- Trace：`outputs/qwen35-swebench-50sess.jsonl`（SWE-Bench 抽样，4449 reqs / **52 sessions** / 每 session 8-150 turns / time-scale=10 / concurrency=32）
+- 模型：Qwen3.5-35B-A3B (TP4) 和 Qwen3-30B-A3B (TP1) 两组
+- 硬件：单机 8×H100 80GB，mooncake TCP loopback 模拟 P→D 传输
+
+---
+
+# 第一部份：性能数据现象
+
+## 1.1 三种机制在 Qwen3.5-35B (TP4) SWE 50sess 上的表现
+
+来源：`outputs/swebench-exps/`。
+
+| Run | Mechanism | Policy | Errors | Lat mean | Lat P50 | Lat P99 | TTFT mean | TTFT P50 |
+|---|---|---|---:|---:|---:|---:|---:|---:|
+| `pd-disaggregation-default-20260426T202540Z` | pd-disagg | default | **0/4449** | 1.66s | 0.97s | 7.68s | 0.45s | 0.34s |
+| `pd-colo-default-20260426T210129Z` | pd-colo | default | **4447/4449** | – | – | – | – | – |
+| `pd-colo-default-20260427T033519Z` | pd-colo | default | **0/4449** | 1.77s | 0.86s | 9.67s | 0.29s | 0.25s |
+| `pd-colo-kv-aware-20260427T042034Z` | pd-colo | kv-aware | 469/4449 | 1.52s | 0.82s | 8.27s | 0.26s | 0.23s |
+| `pd-colo-kv-aware-20260427T044944Z` | pd-colo | kv-aware | **0/4449** | **1.57s** | 0.81s | 8.48s | **0.22s** | **0.17s** |
+| `kvcache-centric-default-worker-admission-20260426T210800Z` | KVC | default | **4390/4449** | – | – | – | – | – |
+
+### 现象解读
+
+**(1) pd-disagg 是稳定基线**：1.66s mean / 0 errors / 4199 cache hits（94.4%）。可以正常服务。
+
+**(2) pd-colo（DP）有两次 run，第一次几乎全 crash，第二次稳定**：
+- 04-26 的 4447/4449 errors 来自 SGLang `--disaggregation-mode null` + Qwen3.5-35B-A3B（Mamba/GDN hybrid）的 `token_to_kv_pool_allocator memory leak` bug，crash 了
+- 04-27 的两次 pd-colo run 都跑通了。**`pd-colo-kv-aware-20260427T044944Z` 是这一组实验里跑分最好的配置**——0 errors / TTFT P50 = 0.171s（pd-disagg 的 50%）
+
+**(3) KVC 在 SWE 35B 上的唯一一次 run 几乎全 crash**：4390/4449 = 98.7% errors。但**那 56 个跑通的 direct-to-D 请求性能优异**——Lat mean 1.24s，TTFT P50 0.081s，KV transfer 196 块（vs PD disagg 的 105K 块，**−99.8%**）。说明 KVC 机制本身有效，但 admission control 把绝大多数请求过滤掉了。
+
+### 一句话：在 Qwen3.5-35B 上，**pd-colo + kv-aware 是头名**，KVC 机制配置不当几乎不可用。
+
+---
+
+## 1.2 同 trace 切到 Qwen3-30B (TP1)：v1→v6 演进
+
+为绕开 Mamba 模型的 SGLang bug，团队后续切到 Qwen3-30B-A3B (TP1) 跑 KVC 调优 sweep。**所有结果用同一份 SWE 50sess trace**，可以横向比较。来源：`outputs/qwen3-30b-tp1-*` 各目录。
+
+### 1.2.1 各版本配置概览
+
+| 版本 | 关键改动（一句话） |
+|---|---|
+| v2 | KVC + `--policy default`（这个 policy 选择 **是 bug**，下文 §2.5） |
+| v3 | KVC + `--policy kv-aware` |
+| v4 | v3 + replay 端 session soft_cap 从 4 抬到 16 |
+| v5 (Option D) | 把 admission 决策从 replay 估算改成 D worker 真实容量回答（`worker-mode admission`） |
+| v5+profile | v5 + 1Hz `/server_info` polling 做时序 instrument |
+| v6 P0 | v5 baseline 同配置 rerun ×3 验证可复现性 |
+
+### 1.2.2 各版本同 trace 结果总表
+
+| 版本 | Errors | Lat mean | Lat P50 | Lat P90 | Lat P99 | TTFT P50 | direct-to-D% |
+|---|---:|---:|---:|---:|---:|---:|---:|
+| **8-way DP cache-aware** | **0** | **1.43s** | **0.65s** | **3.61s** | **8.37s** | **0.093s** | – |
+| v3 1P7D KVC | 363 (8.2%) | 4.88s | 1.75s | 12.67s | 28.72s | 0.363s | 39% |
+| v3 2P6D KVC | 9 (0.2%) | 3.58s | 1.52s | 9.23s | 18.70s | 0.328s | 31% |
+| v4 1P7D cap=16 | 435 (10%) | 4.21s | 1.08s | 13.38s | 24.45s | 0.056s | 49% |
+| v4 2P6D cap=16 | 403 (9%) | 2.51s | 0.84s | 6.51s | 18.34s | 0.051s | 53% |
+| v5 1P7D Option D | 9 (0.2%) | 5.18s | 1.59s | 14.67s | 26.09s | 0.207s | 45% |
+| v5 2P6D Option D | 9 (0.2%) | 3.49s | 1.31s | 9.09s | 24.92s | 0.244s | 41% |
+| v5+profile 1P7D | 6 (0.1%) | 4.21s | 1.18s | 11.33s | 28.83s | 0.060s | 55% |
+| v5+profile 2P6D | **415 (9.3%)** | 3.23s | 1.11s | 8.36s | 20.26s | 0.168s | 41% |
+| v5 rerun ×3（无 profile） | **372 / 912 / 396** | 3.00–3.50s | 0.94–1.22s | 7.68–8.65s | 18.97–20.37s | 0.07–0.18s | 40-42% |
+
+**8DP CA 在每一项指标都是头名**：
+- Latency mean **比所有 KVC 配置好 +43%~+260%**
+- TTFT P50 **0.093s**（KVC 最佳 v4 2P6D 是 0.051s——TTFT 单项 KVC 是有优势的，但被整体 P99 灾难抵消）
+- 0 errors（KVC 任一配置 errors 在 9-912 之间漂移）
+
+### 1.2.3 v5+profile 的诡异：加 1Hz polling 让 errors 从 9 涨到 415
+
+这条单独看：v5 baseline 跑出来 9 errors，加上 1Hz `/server_info` polling 之后 415 errors（**46×**）。原因机理见 §2.5。
+
+### 1.2.4 v6 P0 用 ×3 rerun 验证可复现性，结果是不能复现
+
+**关键事实**：v5 baseline 完全相同配置跑 3 次：
+
+| Run | Errors | Lat mean | Lat P50 | TTFT P50 |
+|---|---:|---:|---:|---:|
+| rerun1 | **372** | 3.50s | 1.11s | 0.147s |
+| rerun2 | **912** | 3.00s | 0.94s | 0.071s |
+| rerun3 | **396** | 3.42s | 1.22s | 0.183s |
+
+errors 漂移 **2.5×**（372→912）。Latency mean / P50 也漂移 ~30%。**这意味着 v3-v6 之前所有"single-run"对比的差异 < 30% 的都不可信。**
+
+但要注意：**3 次 v5 中最优的 P50（0.94s）仍然比 8DP CA（0.65s）慢 1.45×**——这个差距大于 single-run variance，所以"DP 全胜 KVC"的头条结论不受 variance 影响。
+
+### 1.2.5 一个有趣的反差：v4 vs v5
+
+- v4：errors 多（~10%）、direct-to-D 占比高（53-58%）、整体 P50 较好（0.84s）
+- v5：errors 少（0.2%）、direct-to-D 占比降低（41-45%）、整体 P50 反而退步（1.31s）
+
+**v5 没有让性能变好，只是把"硬错误"转成了"诚实拒绝"——v4 的 admission 是乐观估算，admit 进来后 D 装不下变成 mooncake 32s timeout（统计成 errors）；v5 让 D 自己拍板，admit 拒得早，请求改走 fallback（统计成低 direct-to-D 率）。容量本身没变。**
+
+---
+
+## 1.3 microbench 上 KVC 击败 PD disagg —— 但本仓库没保留实际 run
+
+`docs/PROJECT_OVERVIEW.md` 写明：
+
+> micro-benchmark 上，`kvcache-centric` 可以比 `pd-disaggregation` 好。原因很简单：**session 少、D KV 放得下**，turn2+ 可以直接走 D session。
+
+但 `outputs/` 里**没有** microbench 实际 run（只有 microbench trace 生成器 `microbench.py` 和它的几个示例 trace 文件）。所以 microbench 的"KVC 赢"是基于设计预期 + 历史口口相传，**没有可重现的产物**。
+
+**这本身是个问题**——下文 §2.6 会解释 microbench 的默认参数（4 sessions × 30K input × 1K append）正好把所有 KVC 失效条件都规避掉了。
+
+---
+
+## 1.4 头条结论（Part 1 总结）
+
+| 工作负载 / 模型 | 头名机制 | KVC 表现 |
+|---|---|---|
+| Microbench（8 session × 30K × 1K append） | KVC > PD disagg（无落地数据，按设计） | 设计上必然赢 |
+| SWE 35B (TP4) | **pd-colo + kv-aware**（1.57s mean, 0 errors） | KVC 唯一 run 中 98.7% errors |
+| SWE 30B (TP1) | **8-way DP cache-aware**（1.43s mean, 0 errors） | KVC 6 个配置全输；最佳的 v4 2P6D 慢 75%、errors 9% |
+
+**真实 agentic 工作负载（SWE-Bench）上，KVC 机制目前没有任何配置能跑赢 naive DP cache-aware。**
+
+---
+
+# 第二部份：结构性问题分析
+
+每条按 (1) 现象（实锤数据）、(2) 根因（代码位置）、(3) 影响量化 三段交代。
+
+## 2.1 KvAwarePolicy 不感知 D 容量 + Session 永久 pin 在初始 D 上 ★ 最严重
+
+### 2.1.1 现象（实锤）
+
+**(a) 每个 session 整 run 中只访问 1 个 D**——基于 v5 rerun1/2/3 全部 4449×3 = 13347 条 metrics：
+
+| Run | sessions | avg distinct-D-per-session |
+|---|---:|---:|
+| rerun1 | 52 | **1.00** |
+| rerun2 | 52 | **1.00** |
+| rerun3 | 52 | **1.00** |
+
+3 次独立 run、156 次 session 实例，**没有一个** session 跨 D 迁移过。
+
+**(b) Direct-to-D 命中率呈极端双峰**——以 rerun1 为例（其他两次形态相同）：
+
+| direct-to-D rate | session 数 |
+|---|---:|
+| 0–20%（"饿死"） | **15** |
+| 20–40% | 7 |
+| 40–60% | 11 |
+| 60–80% | 5 |
+| 80–100%（"顺利"） | **14** |
+
+中间档稀少，两端拥挤。
+
+**(c) 跨 3 次 run 一致饿死的 session = 13/52，且这些 session 的 input 是顺利 session 的 1.98×**：
+
+```
+13 sessions starved (<20% direct-to-D) in ALL 3 runs
+  avg peak input of consistently-starved sessions: 62043 tokens
+  avg peak input of consistently-lucky sessions:   31344 tokens
+```
+
+**结构性、可复现、与 session 大小强相关。** 排除"运气"假说。
+
+### 2.1.2 根因（代码）
+
+`policies.py:166-172` `KvAwarePolicy.select()` 评分函数：
+
+```python
+score = (
+    overlap + sticky * self.sticky_bonus,    # 主项：历史 KV overlap
+    sticky,                                   # 二级
+    inflight_penalty,                         # 三级
+    assignment_penalty,                       # 四级
+)
+```
+
+**评分中完全没有 D 当前容量项**。
+
+session X 第一次落到 D-2 → 在 D-2 上积累 hash_id → 之后不管 D-2 多满，X 的 turn N+1 的 overlap 在 D-2 上仍是最大 → 永远选 D-2。即使 D-5 全空也轮不到。
+
+`RoutingState.decode_resident_blocks` (`policies.py:46`) 还从不缩减——但因为 SWE trace 的 hash_ids 是 session-unique，**不缩减并不影响"选对 D"，只影响内存**——真正问题在评分函数无容量项。
+
+### 2.1.3 影响量化
+
+- 25%（13/52）的 session 几乎每个 turn 走 fallback 路径
+- fallback 路径 mean lat 约 3.5s vs direct-to-D ~0.5s——**饿死 session 每 turn 慢 6×**
+- 这 13 个 session 还容易撞 mooncake 32s timeout（见 §2.2、§2.3），P99 完全由它们决定
+- **SLO 视角下：25% 的用户体验是系统性糟糕**
+
+---
+
+## 2.2 D 端 LRU 只能 evict idle session → 跟不上压力
+
+### 2.2.1 现象（实锤）
+
+来源：`outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/.../logs/decode-{0..5}.log`，全 run 计数：
+
+| D worker | "Trimmed decode session cache" 事件 | KVTransferError | 峰值 token_usage |
+|---|---:|---:|---:|
+| decode-0 | 9 | 0 | 0.99 |
+| decode-1 | 43 | 4 | 0.99 |
+| decode-2 | 16 | **153** | 0.97 |
+| decode-3 | 37 | 29 | 0.99 |
+| decode-4 | 28 | **90** | **1.00** |
+| decode-5 | 30 | **93** | **1.00** |
+
+**所有 6 个 D 都顶到 token_usage ≥ 0.97，2 个顶到 1.00（KV 池完全耗尽）。LRU 触发 9-43 次，远不够——transfer 错误是 LRU 触发量的 5-10×。**
+
+decode-2 极端：trim 16 次 vs error 153 次 = LRU 跑得比错误慢 9.5×。
+
+### 2.2.2 根因（代码）
+
+`scheduler.py:2040` 的 `evict_idle_streaming_sessions_lru` 实际只能 evict：
+
+> 所有 req 都 finished + streaming 模式 + 该 session 没有 inflight transfer
+
+但 SWE 高并发（concurrency=32 + time-scale=10 → effective inter-turn gap p50=0.25s）下，每个 session 几乎一直有 inflight req。**hot session 永远不 idle，LRU 永远找不到东西可踢。**
+
+### 2.2.3 影响量化
+
+- 单 run 累计 KVTransferError：6 个 D 之和 = **369 次**
+- 对应 ~8% 请求失败率（v5 errors 9/372/912 三次平均 ~430/4449 = 9.7%）
+- **每次 mooncake timeout = 32s**——直接构成 P99 18-26s 的尾巴
+
+修复需要 SGLang 内部分层 eviction：除 idle session 外，按访问频率 / 时序加权强制 retract——**不在当前 KISS 边界**。
+
+---
+
+## 2.3 没有 D → Replay backpressure 通道
+
+### 2.3.1 现象
+
+§2.2 数据显示 D 顶到 token_usage=1.00 时仍在持续接收新请求，最终撞 mooncake 32s timeout。**整个错误链路里没有"D 过载，请慢点发"的反向信号**。
+
+定量证据：rerun1 的 KVTransferError 时间分布——**98% 集中在 run 后半段**（参考 `KVC_DEBUG_JOURNEY_V1_TO_V5.md` §v4）。前期 D 容量充裕时正常，达到上限后**所有后续请求集中失败**——典型的"无 backpressure 系统在过载点雪崩"模式。
+
+### 2.3.2 根因（代码）
+
+链路：
+
+```
+replay 端按 trace 时序 + concurrency=32 持续发请求
+  ↓
+PD Router 裸 round-robin (pd_router.py:43-49)
+  ↓
+P 收到请求做 prefill → mooncake 推 KV → D 端
+  ↓
+D 端 transfer queue 堆积 → 32s timeout
+  ↓
+errno 抛回 replay → fallback 路径，但 concurrency 不降
+```
+
+D 端的 `admit_direct_append` 响应里**只有 can_admit/reason 等过去时字段，没有任何"建议节流"的指示**。
+
+### 2.3.3 修复（本次代码改动已实现）
+
+代码已加 `recommended_pause_ms` 字段：
+- `third_party/sglang/.../io_struct.py:DirectAppendAdmissionReqOutput` 增加 `recommended_pause_ms: int = 0`
+- `scheduler.py:_compute_backpressure_pause_hint`：按 `transfer_queue_depth`、`retracted_queue_depth`、`token_usage_after` 计算
+- `replay.py`：admission 响应里读到 hint → 更新 `DecodeResidencyState.pause_until_s[D]` → 下次发到该 D 之前 sleep
+- CLI flag：`--enable-backpressure`（默认 off，保留 baseline 行为）
+- 同时新增 3 个结构性日志（`structural/admission-events.jsonl` / `backpressure-events.jsonl` / `session-d-binding.jsonl`）
+
+**待 GPU smoke 验证。预期 errors 从 ~370 降到 < 50；P99 改善（消除 32s timeout 尾巴）；mean latency 可能略升（被强制 sleep）。**
+
+修复脚本：`scripts/sweep_backpressure_smoke.sh`（4 个 run × 30-60 min）；分析器：`scripts/analysis/analyze_backpressure_smoke.py`。
+
+### 2.3.4 注意
+
+backpressure 是**降级机制**，不是性能优化——它把"硬错误（32s timeout）"换成"主动等待"。整体 throughput 不会因此提升，但 P99 应大幅改善。
+
+---
+
+## 2.4 P-side round-robin 不感知 D 健康
+
+### 2.4.1 现象（实锤）
+
+来源：v5 rerun1 `prefill-{0,1}.log`，全 run 计数：
+
+| Worker | KVTransferError | "Decode instance could be dead" | 请求量 |
+|---|---:|---:|---:|
+| prefill-0 | **367** | 361 | 2225 |
+| prefill-1 | **2** | 0 | 2224 |
+
+**两 P 请求量完全均衡（round-robin），错误率差 180×**。日志里 prefill-0 的失败反复指向某个特定 D 的 IP（`to 10.45.80.47:XXXXX`）。
+
+### 2.4.2 根因（代码）
+
+`pd_router.py:43-49`：
+
+```python
+prefill_url, bootstrap_port = self.config.prefill_urls[
+    self.prefill_cursor % len(self.config.prefill_urls)
+]
+self.prefill_cursor += 1
+```
+
+裸 round-robin。不感知：
+- P 当前 inflight transfer 数
+- 目标 D 的健康状态 / 容量
+
+后果：当某个 D 进入 hot 状态时，被 round-robin 派去给它推 KV 的 P **持续失败**；另一个 P 接到的请求恰好命中健康 D，完全没事。**单 P 故障不会被路由层避开。**
+
+### 2.4.3 影响量化
+
+- prefill-0 几乎独自承担了**全部 KVTransferError 的 99%**（367/(367+2)）
+- 如果 router P 选择能避开"正在和 hot D 死磕"的链路，这部分 ~8% 的整体错误率应可降到 < 1%
+
+### 2.4.4 备注
+
+这条结论目前来自单次 run 的 N=1 数据。需要跨 N≥3 次 rerun 验证一致性才能完全确信——加上 §2.1.1 (b/c) 也证明 P-D 链路绑定结构性强相关，"prefill-0 死磕某 D"很可能在每次 run 都重复（由初始 session 落点决定）。
+
+---
+
+## 2.5 Admission RPC 进 scheduler 主循环 → 自我干扰
+
+### 2.5.1 现象（实锤）
+
+v5 baseline 配置不开 polling：errors = 9
+完全相同配置 + 1Hz `/server_info` polling：errors = **415**（**46×**）
+
+来源：`outputs/qwen3-30b-tp1-v5-optD/exp2_2p6d_kvc_optD_summary.json`（baseline 9 errors）vs `qwen3-30b-tp1-v5-optD-profile/exp2_2p6d_kvc_optD_profile_summary.json`（415 errors）。
+
+### 2.5.2 根因（代码）
+
+`/server_info`（被 polling 调用）和 `admit_direct_append` 都进 SGLang scheduler 主循环：
+
+- `/server_info` → `scheduler.py:get_streaming_session_cache_status` → 遍历每个 session slot 计算 `is_idle`
+- `admit_direct_append` → 读 `token_to_kv_pool_allocator.available_size()` + 触发 `maybe_trim_decode_session_cache`
+
+scheduler 主循环本身在跑 decode/prefill 的 forward。这些 RPC 进队列就和 forward 抢调度。
+
+### 2.5.3 真实负载下 admission RPC 频率远高于 1Hz
+
+- 4449 reqs / ~2700s ≈ **1.6 reqs/s**
+- 每个 turn 做 1-3 次 admission probe（direct-append + 可能的 seed retry）
+- × 8 worker = **每秒 ~16-40 次 admission RPC**
+
+也就是 admission 流量本身比 1Hz polling 高一个量级。如果 1Hz polling 都能让 errors 涨 46×，admission 自己的扰动至少同等。
+
+### 2.5.4 修复
+
+不在本轮 KISS 内。设计方向是把 admission 拆成两个端点：
+- `POST /probe` → lock-free 读 snapshot（轻），90% 流量走这条
+- `POST /commit_evict` → 进 scheduler 队列，做实际 LRU（重），仅 probe 不够时调
+
+这部分需要 SGLang 内部 atomic publish snapshot 到共享内存——**结构性改动**。
+
+### 2.5.5 注意
+
+v6 P0 的 ×3 baseline rerun（不开 polling）errors 也是 372/912/396——**polling 不是 415 唯一原因**。本身 v5 admission 设计就敏感，polling 是放大器。
+
+---
+
+## 2.6 Replay 时间被 time-scale=10 压缩 → 测量学失真
+
+### 2.6.1 现象（实锤）
+
+v5 rerun1 metrics 解出的真实 inter-turn gap 分布：
+
+```
+原始 trace inter-turn gap (n=4397):
+  p10=1.6s   p50=2.5s   p90=7.8s   p99=25.1s   max=261s
+
+time-scale=10 实际 replay gap (= 原始 / 10):
+  p10=0.16s  p50=0.25s  p90=0.78s  p99=2.5s    max=26s
+```
+
+### 2.6.2 这意味着什么
+
+真实 agentic 用户/agent 在每个 turn 之间停 **2-8 秒**——思考、打字、tool call 异步返回、agent reasoning。
+
+`microbench.py:20-21` 的默认 `inter_turn_gap_s=1.0` + `session_stagger_s=0.1` 也大致符合这个量级（1 秒左右）。
+
+但 SWE replay 设的 time-scale=10 把这个间隔**人为压到 0.25 秒**——D 还没消化完 turn N，turn N+1 就来了。
+
+### 2.6.3 为什么这么设计
+
+纯粹**节省测试时间**：
+- 原始 trace 跨度 ~6000s（≈100 分钟）
+- time-scale=10 → ~600s（≈10 分钟）
+- sweep 5 版本 × 3 重复 = 25h vs 2.5h
+
+### 2.6.4 它扭曲了什么
+
+1. **抹掉 D 的自然 idle 时间**：真实部署里每个 session 在 turn 间有几秒空窗，正好让 D 端 LRU 把它 evict 出去给其他 session 让位（§2.2 idle 判定）。time-scale=10 下几乎所有 session 一直忙——LRU 永远找不到 idle session。
+2. **人为提升并发压力**：concurrency=32 在 time-scale=10 下意味着 D 端持续承受 320 effective concurrent agents 的压力——远超真实部署。
+3. **掩盖 backpressure 等慢节奏机制的价值**：如果 inter-turn gap 是 2.5s，backpressure 让 replay 等 0.5s 几乎不影响吞吐；time-scale=10 下 0.5s 的 sleep 等于直接跳过下一个 turn。
+
+### 2.6.5 严重性：所有 KVC vs DP 结论都带这个失真
+
+**v3-v6 全部数据基于 time-scale=10**。所以"KVC 在 SWE 上输给 DP"的程度可能被 benchmark 放大。**真实部署里 inter-turn gap 是 2.5s 的话，KVC 可能根本不会撞到当前看到的容量瓶颈**。
+
+这是项目当前**最严重但还没修的测量学问题**。修复成本极小（只是去掉 `--time-scale 10`），但意义重大——**P0 应该立刻跑一组 time-scale=1 baseline**（KVC + DP 各 N=3）。
+
+---
+
+## 2.7 direct-to-D append 阈值 = 2048 是个 magic number
+
+### 2.7.1 现象（实锤）
+
+`replay.py:51` 默认值：
+
+```python
+kvcache_direct_max_uncached_tokens: int = 2048
+```
+
+判定（`replay.py:2177`）：当新 turn 的 uncached append > 2048 token 时，**禁止 direct-to-D**，请求改走 P→D reseed 路径。
+
+实测 v5 rerun1 的 uncached append 分布（`input_length - cached_tokens`）：
+
+```
+所有 4449 请求:
+  p10=50  p25=181  p50=610  p75=2907  p90=36495  p99=91600  max=103971
+
+> 2048: 1222/4449 = 27.5%
+```
+
+**双峰分布**：median 只有 610，但 p90 已经 36K。
+
+### 2.7.2 根因（代码）
+
+阈值是个 magic number——**没有任何代码注释解释为什么是 2048**，git log 里也没人调过它。
+
+合理推测它存在的理由（按可信度）：
+
+| 理由 | 是否成立 |
+|---|---|
+| D 是 decode-tuned，max-prefill-tokens 通常 4-8K，append > 2K 会触发 D 内部多 chunk prefill 拖慢 decode | 强 |
+| 大 append 在 D 上 prefill 会阻塞当前正在 decoding 的其他 session 的 TPOT | 强 |
+| P 有更优化的 prefill kernel 和 batch | 弱（D 的 prefill kernel 同源） |
+| 工程上的"安全默认值"，没认真测过 | 强（git log 印证） |
+
+### 2.7.3 但更严重的 bug：execution_mode 标签命名错位
+
+`execution_mode` 名字里带 "large-append" 的请求一共 **2060 个**，其中：
+
+- **1222 个（59.3%）实际 uncached append ≤ 2048**
+
+也就是说，**"large-append" 这个标签名对超过一半的实例是错的**。看 `replay.py:2168-2178` 的判断：
+
+```python
+if (
+    _should_bypass_prefill(...)              # 要求 overlap > 0
+    and direct_append_length is not None
+    and direct_session_reused                 # 要求 session 在本 D 上 opened 过
+    and not direct_session_reset
+    and direct_append_length <= config.kvcache_direct_max_uncached_tokens
+):
+    # direct-to-D
+else:
+    # 进入 "large-append" 分支
+```
+
+**这个 else 分支的 5 个进入条件里，"append > 2048" 只是其中一个。** session 不在本 D 上、被 evict 过、overlap=0 都会进这个分支，但 `execution_mode` 仍然写 `pd-router-fallback-large-append-*`——导致看 metrics 的人误以为问题是 append 太大。
+
+### 2.7.4 实际：阈值不是主要瓶颈，session 不在 D 上才是
+
+把 turn≥2 的请求按"append 是否 > 2048"和"实际 execution mode"交叉：
+
+```
+Turn≥2 小 append (≤2048), n=3129:
+   1854 (59%)  kvcache-direct-to-d-session            ← 走通了
+   1141 (37%)  pd-router-fallback-large-append-session-cap  ← 标签骗人
+   ...
+
+Turn≥2 大 append (>2048), n=1216:
+    813 (67%)  pd-router-fallback-large-append-session-cap
+    365 (30%)  kvcache-centric (失败)
+     22         pd-router-large-append-reseed                ← 真正受阈值影响的
+    ...
+```
+
+**真正因 append > 2048 而失败的请求**：约 50 个（large-append-reseed + 部分 large-append fallback），仅占总数 1-2%。
+
+**绝大多数 fallback 实际是 §2.1 的 session 不在 D 上**——名字里带 "large-append" 是误导。
+
+### 2.7.5 修复
+
+两件事：
+1. 把 `execution_mode` 标签按真实原因细分——把 "large-append" 拆成 "session-not-resident" / "real-large-append" / "session-reset" 等
+2. 阈值本身可以做 sweep（2048 / 4096 / 8192 / 16384）找最优——但收益空间有限（最多改善那 1-2% 的请求）
+
+---
+
+## 2.8 跨 run variance 巨大：N=1 不可信
+
+### 2.8.1 现象（实锤）
+
+v5 baseline 完全相同配置跑 3 次（`qwen3-30b-tp1-v5-optD-baseline-rerun/`）：
+
+| Run | Errors | Lat mean | Lat P50 | TTFT P50 |
+|---|---:|---:|---:|---:|
+| rerun1 | 372 | 3.50s | 1.11s | 0.147s |
+| rerun2 | **912** | 3.00s | 0.94s | 0.071s |
+| rerun3 | 396 | 3.42s | 1.22s | 0.183s |
+
+errors 漂移 **2.5×**（372→912），P50 latency 漂移 ~30%，TTFT P50 漂移 **2.6×**。
+
+### 2.8.2 根因（推测）
+
+源头不止一个，至少包含：
+
+1. **§2.1 + §2.2 的复合**：D 容量过载是临界点附近的非线性系统——initial session-to-D assignment 的随机性决定了哪个 D 先饱和。
+2. **mooncake TCP loopback 的随机性**：单机 loopback 的 32s timeout 触发概率受当前 GPU 内存碎片、PCIe 状态影响。
+3. **scheduler 主循环里 admission RPC 与 decode 抢资源的随机性**（§2.5）。
+
+### 2.8.3 影响
+
+**所有 single-run 比较 < 30% 差异都不可信**。这意味着：
+- v3 vs v4 的 P50 差异（1.75s vs 1.08s）勉强有意义（差异 38%）
+- v4 vs v5 的 P50 差异（0.84s vs 1.31s）勉强有意义（差异 56%）
+- v5+profile 的 1P7D vs baseline（mean 4.21s vs 5.18s）→ 差异 18%，**不可信**
+- 所有 `direct-to-D 占比 ±5%` 的差异都是噪声
+
+### 2.8.4 这条规则要求所有后续实验
+
+**要任何 KVC 配置间或 KVC vs DP 的对比，最少跑 N=3，最好 N=5。** 不跑 N≥3 的实验在做"碰运气科研"。
+
+8h 一次 sweep 装不下 N=3 + 多版本对比，所以必须**牺牲版本数量保 N≥3**。
+
+---
+
+## 2.9 microbench 的 KVC 优势不能外推到真实 agentic
+
+`microbench.py:13-22` 默认参数：
+
+| 维度 | 默认值 |
+|---|---|
+| `session_count` | 8 |
+| `turns_per_session` | 3 |
+| `initial_input_length` | 10000 |
+| `append_input_length` | **1000** ← 低于 §2.7 的 2048 阈值 |
+| `output_length` | 1000 |
+| `inter_turn_gap_s` | **1.0** ← 接近真实 agentic |
+| `session_stagger_s` | 0.1 |
+
+**与 SWE workload 的关键维度对比**：
+
+| 维度 | microbench | SWE 50sess |
+|---|---|---|
+| Session 数 | 4-8 | 52 |
+| Per-session peak input | ~31K | median 49K, max 104K |
+| 总 working-set / 7D 容量（92K each） | 0.19×（5× 冗余） | **3.95×（4× 过载）** |
+| Append size 是否过 2048 | 几乎 100% 过不到 | 28% 超过 |
+| Session 数是否过 cap | 4 ≤ 28（v3 cap×7D） | 52 远超 |
+
+**Microbench 把 KVC 的所有失效条件都规避了**：容量充裕、append 卡阈值之下、session 数远低于 cap、inter-turn gap 接近真实——这一组参数让 KVC 五项判断（路由 / admission / 没被 evict / append ≤ 阈值 / 无 backpressure）全部通过 → 100% 走 direct-to-D 快路径。
+
+**而 SWE workload 在每一项上都把 KVC 推过临界点。**
+
+所以"KVC 在 microbench 赢 PD disagg"是个**弱命题**——它只证明了机制能跑，没有证明在真实 agentic 下能赢。
+
+---
+
+# 第三部份：一句话总结与下一步
+
+## 现状一句话
+
+> 在所有可比的真实 agentic workload（SWE 35B / 30B）上，**naive DP cache-aware 全胜 KVC 任何配置**，且差距 > 30%（远超 single-run variance）。Microbench 上 KVC 赢 PD disagg 的设计前提（容量富余、append 小、session 少）在真实 workload 下不成立。
+
+## 排序后的结构性问题（按修复 ROI）
+
+| 排名 | 问题 | 影响 | 修复成本 |
+|---|---|---|---|
+| **P0** | §2.6 time-scale=10 失真 → 所有 KVC vs DP 结论可能被 benchmark 放大 | 颠覆性 | 极低（改 flag） |
+| **P0** | §2.1 session 永久 pin + 容量盲选 | 25% session 永远饿死 | 中（改 policy） |
+| **P0** | §2.2 D-side LRU 跟不上 | ~8% errors 来自此 | 中（改 SGLang） |
+| P1 | §2.3 没 backpressure | 把 timeout 雪崩变可控 | **已实现**（待 GPU smoke） |
+| P1 | §2.4 P-side 不感知 D 健康 | 单 P 出错率差 180× | 中 |
+| P1 | §2.7 / 2.8 metrics 标签命名错位 | 数据解读经常出错 | 低（改字符串） |
+| P2 | §2.5 admission RPC 进 scheduler 主循环 | 自我干扰 | 高（结构改动） |
+| P2 | §2.8 N=1 不可信 | 实验方法学 | 0（团队约定） |
+
+## 立刻能做的三件事
+
+1. **跑 time-scale=1 baseline**（KVC v5 + 8DP CA 各 N=3，~6h GPU）—— 不修代码、单变量、决定后续路线。
+2. **跑 backpressure smoke**（已实现，4 run × ~30-60 min，~3-4h GPU）—— 验证 §2.3 修复的端到端效果。
+3. **修 metrics 标签命名**（`pd-router-fallback-large-append-*` → 按真实原因分类）—— 让以后看数据的人不会再被误导。
+
+## 不立刻做但要重新讨论的
+
+- **§2.1 capacity-aware policy**：之前考虑过的"评分加 capacity 项"会引入"换 D"的副作用（孤儿 KV、新 D 上仍可能饿死），需要跟 §2.2 的 D 端 hot retract 一起设计。
+- **§2.5 admission API 拆 probe / commit**：是结构性正确方向，但要动 SGLang 内部 + atomic publish 机制，不是 KISS。
+- **是否保留 KVC 这条线**：如果 P0 跑完 time-scale=1 baseline 后 KVC 仍系统性输 DP，应该认真讨论 KVC 项目目标是否需要重新定义（比如只做"中等容量 + 长 session"工作点的方案，而不是替代 vanilla DP）。
+
+---
+
+## 附录 A：本报告所有数据的来源
+
+| 章节 | 数据源 |
+|---|---|
+| 1.1 SWE 35B | `outputs/swebench-exps/{pd-disagg,pd-colo,kvcache-centric}-*` |
+| 1.2 TP1 series | `outputs/qwen3-30b-tp1-{exps,v3-kvaware,v4-cap16,v5-optD,v5-optD-profile,v5-optD-baseline-rerun}/` |
+| 2.1 session pinning | `outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/exp2_2p6d_run{1,2,3}_metrics.jsonl` |
+| 2.2 D LRU 计数 | `outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/.../logs/decode-{0..5}.log` |
+| 2.4 P imbalance | `outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/.../logs/prefill-{0,1}.log` |
+| 2.5 polling 影响 | v5 baseline summary vs v5+profile summary |
+| 2.6 inter-turn gap | rerun1 metrics 的 `trace_timestamp_s` 字段 |
+| 2.7 append 分布 | rerun1 metrics 的 `input_length - cached_tokens` |
+| 2.8 variance | rerun1/2/3 三组 summary |
+
+## 附录 B：相关已有文档
+
+- `docs/PROJECT_OVERVIEW.md` — 项目目标、microbench 结论
+- `docs/AGENTIC_FIT_ANALYSIS_ZH.md` — 结构性缺陷的早期分析（本报告 §2 的来源）
+- `docs/KVC_DEBUG_JOURNEY_V1_TO_V5.md` — v1→v5 详细演进日记
+- `docs/V5_PROFILE_INVESTIGATION_ZH.md` — v5+profile 调查（含 critic 修订）
+- `docs/SWEBENCH_EXPERIMENT_RESULTS.md` — SWE 35B 早期实验
+- `docs/REFACTOR_PLAN_ZH.md` — 当前重构计划
+- `docs/STRUCTURAL_VALIDATION_REPORT_ZH.md` — 结构性 claim 验证（本报告的精简版）

docs/V2_DEEP_ANALYSIS_ZH.md

@@ -0,0 +1,624 @@
+# KVC v2 深度分析：相对 TEAM_REPORT 基线的改进、性能、新暴露的问题
+
+**日期**：2026-05-11
+**对象**：项目团队同学
+**基线**：`docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md`（v3-v6 ts=10 调优 sweep 的状态报告）
+**新数据**：
+- `docs/REFACTOR_PLAN_V1_ZH.md`（ts=1 4-run validation 结果）
+- `docs/MIGRATION_V1_FINDINGS_ZH.md`（v1 thrashing 诊断）
+- `docs/V2_RESULTS_ZH.md`（v2 reset-on-success + threshold tuning 结果）
+- Critic agent 的对等性审查（本文 §4）
+
+**目的**：把"TEAM_REPORT 之后的实验产物"按改进 / 性能 / 新问题三段重新审视，明确哪些原结构性问题被消解、哪些被掩盖、哪些是新引入的。
+
+---
+
+## 0. TL;DR
+
+1. **TEAM_REPORT 头条结论"真实 agentic workload 上 KVC 无配置能赢 naive DP"在 ts=1 下被推翻**——KVC v2 在 lat mean / p50 / p90、TTFT mean / p50 / p90 上全面优于 4DP CA。
+2. **生产决策结论：online coding agent serving 应选 KVC 1P3D**。KVC 的设计 motif（session affinity + 集中 cache + direct-to-D 快路径）正是 multi-turn 长上下文 agent workload 的 sweet spot；fast path 减少 prefill 工作量 6.9× 是机制目标实现，不是 measurement artifact。
+3. **真实代价只有一项：TTFT p99 = 1.29s vs DP 0.43s（KVC 3× 差）**——来自 8.3% 非 direct-to-D 路径的 mooncake reseed 长尾。生产部署要么用真 RDMA 把这条压下来，要么靠容量规划让 reseed 极少发生。
+4. **TEAM_REPORT §1（session pin 饿死）已被 v2 修好**——direct-to-D 从 42.8% 涨到 91.6%，severe thrashing 清零。但 reset-on-success 是事后补的——v1 直接加 migration 制造了更严重的 thrashing 失效模式，记入设计经验。
+5. **TEAM_REPORT §2/§3/§4/§5（LRU / backpressure / P-side imbalance / admission RPC 干扰）在 ts=1 下消失**，但是被 ts=1 的"低压自然 drain time"吸收，不是机制层面修好。一旦回到 ts=10 / 更长 trace / 更紧容量，会全部复现——属于潜在的，不是消除的。
+6. **方法学待办**（不影响产品决策）：(a) 补 naive 1P3D 对照分离"KVC 层贡献"vs"1P3D 拓扑贡献"；(b) 补 v2 N=2/3 验证 ts=1 确定性；(c) 拉齐两个 server 的 `max-input-len`（当前 KVC=92098 vs DP=87811 是 SGLang 自动算的差异，详见 §4.3）。
+
+---
+
+## 1. 三组新实验与 TEAM_REPORT 的关系
+
+### 1.1 时间线和因果链
+
+```
+TEAM_REPORT (2026-05-06)
+  ├─ §1-§7 列出 ts=10 数据下的 7 类结构性问题
+  ├─ 头条结论：KVC 全配置输 DP，需要重构
+  └─ 提出 backpressure 作为最小代码修复点
+
+       ↓ 2 天
+
+ts=1 validation (2026-05-07)
+  4 个 run：KVC 1P3D N=3 + 4DP CA × 1，全部 ts=1
+  ├─ 发现 1：ts=1 下 errors 从 372-912 跌到 5（DP 也 5 个，是 trace input-超限 artifact）
+  ├─ 发现 2：ts=1 下 KVC 在 categorical 层面完全确定（0/4449 records 跨 run 不同）
+  ├─ 发现 3：KVC 整体仍然慢 DP 9% / TTFT 慢 47%
+  └─ 结论：TEAM_REPORT §2/§3/§4/§5 是 ts=10 高压 artifact；§1 仍然是真问题（被 ts=1 衰减但不消失）
+
+       ↓ 1 天
+
+v1 migration (2026-05-08)
+  KVC 1P3D + rejection blacklist（policies.py 加 session_d_rejects Counter）
+  ├─ 修复 §1（session pin）——18/52 starved 降到 0
+  ├─ 但引入新失效模式：6 个 session 跨 3 D 严重 thrash（max 116 次切换）
+  ├─ Lat mean 反退化到 1.758s，TTFT mean 涨到 0.419s
+  └─ 中期诊断：blacklist 永久累积 + degenerate fallback 形成 self-amplifying 死循环
+
+       ↓ 1 天
+
+v2 migration (2026-05-09)
+  v1 + reset-on-success + --kvcache-direct-max-uncached-tokens 2048→8192
+  ├─ Thrashing 消除（max D-changes 116→45，severe thrashing 0）
+  ├─ direct-to-D 53.3%→91.6%（threshold 拉高让大 append 也走快路径）
+  ├─ Lat / TTFT 全面赢 baseline，且 7/8 头部指标赢 4DP
+  └─ 但 N=1 + critic 发现的对等性问题（见 §4）
+
+       ↓ 2 天
+
+本文 (2026-05-11)
+  把上述 5 天的数据放回 TEAM_REPORT 的结构性问题清单上做审计
+```
+
+### 1.2 同 trace 全部数字总表（按时间）
+
+来源：`outputs/qwen3-30b-tp1-*` 系列各 summary.json。**4449 reqs / 52 sessions / Qwen3-30B-A3B (TP1) / 4×H100 80GB**。
+
+| 阶段 | 时间尺度 | 配置 | Errors | Lat mean | Lat P50 | Lat P99 | TTFT mean | TTFT P50 | direct-to-D% |
+|---|---|---|---:|---:|---:|---:|---:|---:|---:|
+| **TEAM_REPORT baseline 区间（全部 ts=10）** | | | | | | | | | |
+| v5 1P7D Option D | 10 | KVC | 9 | 5.18s | 1.59s | 26.09s | 0.207s | – | 45% |
+| v5 2P6D Option D | 10 | KVC | 9 | 3.49s | 1.31s | 24.92s | 0.244s | – | 41% |
+| v5 rerun1 (重测) | 10 | KVC | **372** | 3.50s | 1.11s | 19.49s | 0.147s | – | ~40% |
+| v5 rerun2 | 10 | KVC | **912** | 3.00s | 0.94s | 20.37s | 0.071s | – | ~40% |
+| v5 rerun3 | 10 | KVC | **396** | 3.42s | 1.22s | 18.97s | 0.183s | – | ~40% |
+| 8-way DP CA | 10 | DP-colo | **0** | **1.43s** | **0.65s** | **8.37s** | **–** | **0.093s** | – |
+| **ts=1 validation 区间** | | | | | | | | | |
+| v0 baseline run1 | 1 | KVC 1P3D | 5 | 1.574s | 0.811s | 8.70s | 0.245s | 0.124s | **42.8%** |
+| v0 baseline run2 | 1 | KVC 1P3D | 5 | 1.573s | 0.809s | 8.74s | 0.243s | 0.120s | 42.8% |
+| v0 baseline run3 | 1 | KVC 1P3D | 5 | 1.574s | 0.812s | 8.76s | 0.243s | 0.123s | 42.8% |
+| 4-way DP CA | 1 | DP-colo | 0 | 1.443s | 0.659s | 8.43s | 0.129s | **0.090s** | – |
+| **Migration 区间** | | | | | | | | | |
+| v1 migration | 1 | KVC 1P3D | 6 | 1.758s | 0.773s | 9.92s | 0.419s | 0.057s | 53.3% |
+| **v2 migration (头条)** | 1 | KVC 1P3D | 5 | **1.432s** | **0.576s** | **8.69s** | **0.098s** | **0.042s** | **91.6%** |
+
+**两组关键对比**：
+
+1. **ts=10 → ts=1（同 KVC 配置）**：Lat mean 5.18s → 1.574s（**3.3× 改善**）；errors 9-912 → 5（**~100× 改善**）；direct-to-D 41% → 42.8%（持平，机制不变）
+2. **v0 → v2（同 ts=1，机制改进）**：Lat mean 1.574s → 1.432s（**9% 改善**）；TTFT mean 0.245s → 0.098s（**60% 改善**）；direct-to-D 42.8% → 91.6%（**+48.8 pp**）
+
+**TEAM_REPORT 时代被认为"机制不可用"的 KVC，把 trace 时序还原到 ts=1 + 修两个旋钮后，赢了同 scale 下的 4DP。**
+
+---
+
+## 2. TEAM_REPORT §1-§9 的逐项更新
+
+按原始优先级排序，每条标注"是否仍是问题 / 被什么消解 / 残留风险"。
+
+### 2.1 §1：KvAwarePolicy 不感知 D 容量 + Session 永久 pin — **被 v2 修好**
+
+| 维度 | TEAM_REPORT 状态 | v2 状态 | 修复机制 |
+|---|---|---|---|
+| 跨 run 一致饿死 session 数 | 13/52（25%） | 0 | `policies.py: session_d_rejects` + `replay.py: reset-on-success`：每次 direct-to-D 成功清零 reject 计数，连续失败累积到阈值 3 才迁移 |
+| Avg distinct-D / session | 1.00 | <2（v2 实测 mean=0.6 D-changes/session） | 同上 |
+| direct-to-D % | 41% | 91.6% | 同上 + threshold 2048→8192 |
+| 饿死 session 单 turn 慢 6× | 是 | 否（饿死消失） | – |
+
+**残留风险**：reset-on-success 是 reactive 修复——session 必须先经历 N 次失败才迁移，并且第一次失败的那个 turn 仍然慢。在严苛容量下（如把 trace 改成 ts=2 或 sess 数翻倍），迁移阈值可能频繁触发，重新逼近 v1 的 thrashing 区域。**未在更紧 workload 上验证。**
+
+### 2.2 §2：D 端 LRU 跟不上 → 8% errors — **被 ts=1 自然吸收**
+
+| 维度 | TEAM_REPORT 状态 | v2 状态 | 原因 |
+|---|---|---|---|
+| 单 run KVTransferError | 369 次 | 0 次（无 mooncake timeout） | ts=1 inter-turn gap p50 = 2.5s 给 D 充分 drain 时间 |
+| D 峰值 token_usage | 6 个 D 全顶到 0.97-1.00 | 偶发 0.97-1.00（burst），常态 0.4-0.85 | 同上 |
+| LRU trim 触发次数 | 9-43（远不够） | 不需要——D 自然回落 | ts=1 工作流 |
+
+**残留风险**：这条**没有机制层面修好**。把 ts 调回 10、或者 session 数从 52 增到 100+、或者 model 切到更大、都会立刻让 D 容量重新顶死，LRU 再次跟不上。**TEAM_REPORT §2 是潜在的，不是消失的。**
+
+### 2.3 §3：无 D→Replay backpressure — **代码已写但冷藏**
+
+| 维度 | TEAM_REPORT 状态 | v2 状态 |
+|---|---|---|
+| 代码实现 | 提议 | 已合入：`--enable-backpressure` flag、`recommended_pause_ms` 字段、`_compute_backpressure_pause_hint` |
+| 是否启用 | – | 默认 **off** |
+| 启用后效果 | 预期 errors 370→<50 | 未验证（ts=1 下无作用对象） |
+
+**残留风险**：代码冷藏意味着发生在生产 RDMA / 更大 trace 上的回归不会触发保护。**如果团队决定项目要支持 ts=10 / 更大 sessions，需要把 backpressure 默认 on 并补 smoke 验证。**
+
+### 2.4 §4：P-side round-robin 不感知 D 健康 — **1P 配置不可测**
+
+v2 是 1P3D，单 P，无从测试 P-side 调度。TEAM_REPORT 数据来自 2P6D 配置。
+
+**残留风险**：未来如果扩到 2P+ 必须重新审查 P 侧调度。**当前数据无法支持也无法反驳。**
+
+### 2.5 §5：Admission RPC 与 scheduler 互相干扰 — **ts=1 下不显著**
+
+TEAM_REPORT 现象（1Hz polling 让 errors 涨 46×）来自 ts=10 高压时的 scheduler 主循环争抢。ts=1 下 D scheduler 大部分时间空闲，RPC 进来不阻塞 batched prefill。
+
+**残留风险**：与 §2 同源——属于 ts=10 高压 artifact。
+
+### 2.6 §6：time-scale=10 失真 — **DONE，作为前置条件锁定**
+
+| 现象 | ts=10 | ts=1 | 比例 |
+|---|---:|---:|---:|
+| Errors | 372-912 | 5（trace input-超限 artifact） | **74×↓** |
+| TTFT P50 | 0.07-0.18s | 0.04s | 4.5×↓ |
+| Per-D spread | ±26% | ±3.8% | 7×↓ |
+| Lat P99 | 18-29s | 8.7s | 2-3×↓ |
+
+**REFACTOR_PLAN_V1 把这条当作所有后续讨论的前置条件——ts=10 数据从此不参与 KVC vs DP 比较。**
+
+### 2.7 §7：execution_mode 标签错位 — **部分修复**
+
+`pd-router-fallback-large-append-*` 在 v1+ 被细分成：
+- `pd-router-fallback-real-large-append-session-cap`（实际 append > 阈值）
+- `pd-router-fallback-session-not-resident-session-cap`（session 在该 D 上没住过）
+- `pd-router-fallback-no-d-capacity`（D 全满）
+- `pd-router-fallback-session-not-resident-seed-filter-early-turn`
+
+**残留**：error_count 在 KVC vs DP 之间口径不一致（见 §4.3），未统一。
+
+### 2.8 §8：N=1 不可信 — **ts=1 下规则改写**
+
+| Trace 区间 | N 要求 |
+|---|---|
+| ts=10 高压 | N≥3（v5 rerun 显示 errors 漂移 2.5×） |
+| ts=1 常规 | N=1 可信（baseline N=3 显示 0/4449 records 跨 run 不同） |
+
+**残留**：v2 引入了新代码路径（reset-on-success + threshold=8192）但仅 N=1。新分支是否仍保持 categorical 确定性**未验证**。这是 critic 标 MINOR 但未关闭的点。
+
+### 2.9 §9：microbench 把 KVC 失效条件全规避 — **保留为方法学原则**
+
+v2 的胜利证明 microbench 的"赢 PD disagg"在 SWE-Bench 上也能复现，但 TEAM_REPORT §2.9 的方法学原则仍然成立——micro-benchmark 应该主动构造能触发 fallback 的 workload。
+
+---
+
+## 3. v2 的真实性能拆解（path-level）
+
+v2 整体跑得快不仅因为 "KVC 机制好"，更因为 **91.6% 请求被路由到了几乎免费的 fast path**。需要看路径级细节才能理解胜利的来源。
+
+### 3.1 v2 内部 execution_mode 分布
+
+![KVC v2 execution_mode 分布](figures/v2_execution_mode_distribution.png)
+
+数据来源：`outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl`，n = 4449（全部请求，含失败）。绿色 = direct-to-D 快路径 = 91.6%；其余红色 = 慢路径 / fallback / 失败。绘图脚本：`scripts/analysis/plot_v2_path_breakdown.py`。
+
+### 3.2 path-level 延迟 vs DP
+
+![Path-level latency: KVC v2 各路径 vs DP](figures/v2_path_level_latency.png)
+
+数据来源：同上 + `outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl`。Y 轴 log 刻度（latency 跨度 41ms ~ 7.71s）。已过滤 abort / error 请求，所有数字按对等口径计算。
+
+**关键事实**：
+- KVC 的 91.6% **fast path** 在 TTFT p50 上是 **41ms vs DP 92ms**——压制 DP 2.2×；TTFT p99 150ms vs DP 428ms 仍优 2.9×
+- KVC 的 **3.4% reseed 慢路径** TTFT p99 = **5.12s**，是 DP 单一路径 p99（428ms）的 **12×**
+- KVC 的 **0.7% no-d-capacity fallback** 是最坏情况：TTFT p99 = 7.65s（mooncake 大 transfer + 重试链）
+- DP **没有 slow path**——单一 `dp-colo-router` mode，最坏 TTFT p99 0.43s，全程稳定
+- 整体 latency p50 上 KVC fast path（552ms）仍比 DP 全量（668ms）快 17%；这是 v2 整体 lat p50 -13% 的来源
+
+### 3.3 Fast path 的工作量比 DP 少 6.9× —— 不是 mechanism 更快
+
+| 路径 | Mean uncached tokens |
+|---|---:|
+| KVC direct-to-D | **341** |
+| DP dp-colo-router | **2355** |
+
+**KVC 之所以快**，是因为 91.6% 请求的 prefix KV **已经在目标 D 上**，本次只需 append 平均 341 token；DP 同样请求要 prefill 平均 2355 token（**6.9× 工作量**）。
+
+这是结构性的 KVC vs DP 差异——**KVC 的设计就是利用 session 间 KV 复用**，所以"工作量少"本身就是机制核心目标。但在比较时必须诚实：
+
+> KVC 的 TTFT 优势 = **session-aware 路由减少了 prefill 工作量**，**不是** D 端硬件层面更快。
+
+如果工作量做归一化（比如限定都做 2000 token 以上 uncached prefill），KVC 应该和 DP 在同一速度量级。
+
+### 3.4 TTFT 概率密度对比：bimodal vs unimodal
+
+把 path-level 数据投影到 TTFT 的分布维度，可以更直观看出 KVC 与 DP 是**本质不同的两种分布形状**：
+
+![TTFT probability density: KVC v2 vs 4-way DP](figures/ttft_pdf_comparison.png)
+
+左图（线性 x ∈ [0, 0.6s]）看 body：
+- **KVC 的 PDF 在 ~40ms 有一个尖锐峰值**（来自 91.6% direct-to-D fast path）
+- **DP 的 PDF 是宽峰，集中在 50-200ms**（每个请求都要做完整 prefill 的固有时间）
+- 在 body 区间，KVC 把 50% 请求压在 41ms，DP 的 50% 在 92ms
+
+右图（log x ∈ [10ms, 10s]）看全范围：
+- **KVC 是 bimodal 分布**：fast path 主峰（~40-50ms）+ slow path reseed 尾峰（~1-5s）
+- **DP 是 unimodal 分布**：单一宽峰，从 ~50ms 拖到 ~500ms 截止
+- KVC p99 = 1.28s 来自小尾峰；DP p99 = 0.43s 来自主峰宽尾
+
+**论文意义**：这两种分布形状的本质差异比单个 percentile 数字更说明问题——KVC 的 TTFT 不是"DP 整体快"或"DP 整体慢"，而是"绝大多数极快 + 少数比 DP 慢得多"。生产决策的判据应该是 **fast path 集中度 vs slow path tail 长度**的权衡，而不是单个 mean 或 p50 数字。
+
+绘图脚本：`scripts/analysis/plot_ttft_pdf.py`（用 `scipy.stats.gaussian_kde`，body 用 Scott bandwidth 0.15，full range 用 log10 域 KDE）。
+
+---
+
+## 4. 需要诚实交代的 caveats（不是 KVC 的设计缺陷）
+
+Critic agent 对 v2 vs 4DP 的对等性做了 10 项审查。下面分两类：
+- **真实代价**（§4.1-§4.3）— KVC 机制本身的开销，无法回避，论文里必须讲清楚
+- **辩驳 critic**（§4.4-§4.5）— critic 把 KVC 的**设计意图**误标为"对比不公平"，本节澄清
+- **方法学待办**（§4.6-§4.7）— 实验对照层面的事，需要补但不影响产品决策
+
+### 4.1 TTFT p99 长尾 — **真实代价，必须显式报告**
+
+实测 TTFT 全分位数：
+
+| 指标 | KVC v2 | DP | Ratio |
+|---|---:|---:|---:|
+| TTFT p50 | 0.042s | 0.090s | 0.47× (KVC 优) |
+| TTFT p90 | 0.091s | 0.252s | 0.36× (KVC 优) |
+| **TTFT p99** | **1.285s** | **0.427s** | **3.01× (DP 劣)** |
+| **TTFT p99.5** | **2.65s** | **0.485s** | **5.47× (DP 劣)** |
+| **TTFT > 1s 计数** | **59** | **9** | **6.5× (DP 劣)** |
+
+之前 `V2_RESULTS_ZH.md §2` 的 headline 表省略了 TTFT p99，是错的。**论文里 headline 必须包含 p99**——KVC 在 mean/p50/p90 全胜但 p99 输 3×，要诚实摆出来。这不是赢负翻盘（p99 之外都赢），但 p99 长尾是真实代价。
+
+### 4.2 TTFT p99 恶化的根因：8.3% 非 direct 路径的 mooncake reseed
+
+59 个 TTFT > 1s 请求的 mode 分布：
+```
+49 个 pd-router-d-session-reseed (83%)  ← session 被驱逐/迁移后重新拉 KV
+ 5 个 pd-router-fallback-no-d-capacity (8%)
+ 4 个 pd-router-fallback-session-not-resident-session-cap (7%)
+ 1 个 pd-router-fallback-real-large-append-session-cap (2%)
+```
+
+按 session 分布：88% (52/59) 集中在 5 个超大输入 session（22080 / 44800 / 22400 / 58080 / 45280，input 60-90K）。
+
+**机理拆分**：reseed 路径的延迟由两段组成——
+1. **P 端 re-prefill 段**：用 trace 中带的完整 prompt 在 P 上重新算 prefill。**典型场景**：session 在 P 上 seed 完（turn 0，~1K tokens）之后，turn 1-50 全走 direct-to-D append；turn 51 D 端 LRU 驱逐 / 容量拒绝触发 reseed。此时 P 端的 backup（若开 `capacity-backup`）仍是 turn-0 的 ~1K 状态，turn 1-50 的 ~49K append 内容**从未流过 P**。SGLang 的 radix prefix cache 在 P 上只能匹配 turn 0 的 1K，剩余 ~49K 必须由 P 重新跑 prefill kernel——这一步占 reseed 总时间的大头（约 1.5-3s @ 1×H100，30B 模型）。
+2. **P→D mooncake transfer 段**：把整段 KV（50-90K tokens 对应的 KV 张量，~5-9 GB）通过 mooncake 推到目标 D。本次 benchmark 用的是 TCP loopback，实测 1.5-4s（取决于 session 大小）。生产用 IB RDMA（节点实际有 mlx5_0/_1 @ 200 Gb/s × 2 active）应可压到 200-400ms。
+
+**两段相加**：当前 reseed 中位 ~2.5s、p99 ~7.7s。
+
+### 缓解策略的真实效果
+
+- (a) **真 RDMA 替换 mooncake TCP loopback**——救的是 transfer 段（~1.5-4s → ~200-400ms），不动 re-prefill 段。预期 reseed 总延迟从 3-7s 压到 **1.7-3.2s**，TTFT p99 从 1.28s 降到 ~0.7s 量级（**仍输 DP 0.43s**）。**当前 sweep 未启用**（缺 `--force-rdma --ib-device mlx5_0`）。
+- (b) **容量规划**：sessions × peak context ≤ 总 D KV pool × 0.7，让 LRU/reseed 几乎不触发。对生产部署而言最可靠，但对本 trace 不适用——sessions 已固定。
+- (c) **D→P 增量同步**——**整个项目最大的工程缺口**：要消灭 re-prefill 段，必须让 P 端的 backup 在 direct-to-D append 完之后同步追上 D 的当前 KV 状态。这样 reseed 时 P 端已经有最新整段 KV，可以直接 P→D transfer，无需 re-prefill。**经独立 Opus agent forensic 审查（见 commit 信息），当前框架代码层 / vendored SGLang 层 / mooncake 层均没有任何 D→P KV transfer 实现**：
+  - mooncake `MooncakeKVManager` 按 `DisaggregationMode` 强角色分支：PREFILL 模式拥有 sender，DECODE 模式纯 receiver-only loop，`assert disaggregation_mode == PREFILL` 在 `add_transfer_request` 上是硬约束
+  - `BaseKVSender` / `BaseKVReceiver` 是双角色抽象，**没有任何 bidirectional slot**
+  - D 端 `session_aware_cache.release_session` 只调 `kv_pool_allocator.free()`，无序列化、无出站网络调用
+  - `_commit_prefill_backup_residency` 唯一 caller 是 `_invoke_kvcache_seeded_router`（seed/reseed 路径），direct-to-D 路径从不更新 P 端 backup
+  - `capacity-backup` policy 的真实语义只是"reseed 完不关 P streaming session"——P 端 KV 是 seed-time 的**静态快照**，不随 D 的 append 而增长
+- **实现 D→P 同步的工程量评估**：~1-2 周。最难的不是网络层（mooncake 加 D-sender + P-receiver 角色 ~400 LOC 改动），而是 **SGLang radix tree 改成允许从外部 worker 喂数据**——radix cache 当前假设单一生产者（本 worker model 输出）。这是论文里 §future-work 的核心 contribution 缺口。
+
+### 4.3 Error 统计口径已修复；abort 数双方都比之前发现的多
+
+之前 V2_RESULTS_ZH.md 说"DP 同样有 5 个 input-too-long abort"。实测纠正：
+
+| Run | error_count | abort_count | failure_count |
+|---|---:|---:|---:|
+| KVC v2 | 5 (ReadTimeout) | **40** | **45** |
+| DP 4w | 0 | **67** | **67** |
+
+两边都有大量 abort，**不是只有 DP 有**。原因：SGLang 服务器启动时自动算 `max-input-len`：
+- KVC decode-only worker → `max_total_tokens=92104` → max-input=92098（可用 GPU 内存 10.85 GB）
+- DP fused worker → `max_total_tokens=87817` → max-input=87811（可用 GPU 内存 8.93 GB，因为还要给 chunked-prefill workspace ~2 GB）
+
+DP 限制更紧，所以 abort 多 27 个。**这是 SGLang 自动 mem 分配的产物，不是机制差异。**
+
+**已修代码**：`src/agentic_pd_hybrid/metrics.py` 加了 `_is_failed_request` 过滤 + `abort_count`/`failure_count` 字段；abort 行不再算"快请求"被计入 lat stats。重算后：
+
+```
+                修复前              修复后（排除 abort）
+KVC v2 lat_mean   1.4323            1.4441
+DP 4w  lat_mean   1.4435            1.4642
+delta (KVC vs DP) -0.8%             -1.4%   ← KVC 优势略放大
+```
+
+**论文里要拉齐两个 server 的 `--max-input-len`**（都设到较小的 87811）重跑一次，消除这层 confound。
+
+### 4.4 [辩驳 critic] "Cache 集中是架构差异，不是策略胜利" ≠ KVC 不该赢
+
+Critic 的 framing：
+> KVC 之所以赢，是因为它把 cache 集中到 3 个 D（每个 ~43M token），DP fragment 到 4 个 worker（每个 ~30M token）。两边 policy 都是 `kv-aware`，差异来自架构而非策略。
+
+**反驳**：KVC 整套机制的**核心设计就是主动选择 affinity 集中而非 fragment**。"差异来自架构"等价于"差异来自 KVC 是 KVC"——这正是要论证的设计点。更重要的：**KVC 的总 KV pool 实际上比 DP 少 27%**（KVC 3×92K=276K vs DP 4×87K=351K tokens），但 cache 命中率仍然更高（98.1% vs 96.8%）。
+
+![Cache efficiency paradox: KVC 用更少的总池子缓存更多](figures/cache_efficiency.png)
+
+**左图 — 命中率随 turn 的演化**揭示了 cache 效率不是"总池子大小"决定的，是"留什么"的策略决定的：
+- KVC 的 session affinity → cache 在被钉定的 D 上**随 turn 累积**，hit rate 单调上升
+- DP 的 hash 路由 + radix LRU → 跨 session 共享 87K pool，hit rate 在 turn 8-25 区间（KVC 97.0% vs DP 95.8%，差 **1.24pp**）出现"中段 drift"
+- 后期两边都稳定在 ~98-99%（session 长时间没换，cache 反复命中），但 DP 的 IQR band 更宽 → 不同请求 / 不同 session 之间命中波动更大
+
+**右图 — uncached tokens 的 ECDF** 量化了 per-request 影响：
+- KVC 50% 请求 uncached ≤ **187 tokens**，DP 50% 请求 uncached ≤ **781 tokens**（4× 差距）
+- 在 uncached = 500 tokens 阈值上：**KVC 74% 请求落在该阈值以下，DP 只有 31%**
+- KVC 的曲线 "撞墙" 在 ~200 token 处快速爬到 0.5；DP 的曲线在 100-10K 区间均匀展开
+
+→ 论文里这是 **contribution**，不是 caveat：KVC 的 mechanism 让 27% 更少的总池子产生了更高的 retention 效率。
+
+### 4.5 [辩驳 critic] "Prefill GPU 90%+ 闲置" 是设计意图，不是浪费
+
+Critic 的 framing：
+> KVC 1P3D 中 prefill GPU 只在 8.3% 请求时被激活；实际工作 GPU 只有 ~3.08 个，对比 4DP CA 的 4 个 fused GPU 不公平。
+
+**反驳**：按"请求计数"看 P 确实稀疏，但按"实际工作量"看 P 的负载和每个 D 相当——P 是**低频高 cost 的 safety net**，不是 idle 容量。
+
+![Per-GPU utilization: 请求计数视图 vs 工作量视图](figures/gpu_utilization.png)
+
+**左图 — 请求计数视图**：KVC P GPU 仅处理 328 个请求（7.4%），而 KVC D 各处理 ~1450 个（33%），DP 各处理 ~1100 个（25%）。**乍看像 critic 说的"P 闲着"**。
+
+**右图 — 工作量视图（compute tokens）**：
+- KVC P GPU：**1.07M tokens 的 prefill 工作**（仅 prefill，无 decode）
+- KVC D GPU 每个：~0.80M tokens（小量 append-prefill + 全部 decode）
+- DP 每个 worker：~1.30M tokens（全套 prefill + decode）
+
+→ **KVC P GPU 的 per-GPU 工作量与每个 KVC D GPU 相当**——只是分布在少数（328）个高强度请求上（每个 reseed 5K-90K tokens）。它不是空转，是 **low-frequency, high-cost safety net**。
+
+**总工作量对比**：
+- KVC 4 个 GPU 合计 ~3.47M tokens 工作
+- DP 4 个 GPU 合计 ~5.17M tokens 工作（**KVC 减少 33% compute**——这是 session affinity 带来的 cache 复用收益）
+
+这两点综合：KVC 用 **同样 4 个 GPU、更少总 KV pool、更少总 compute**，做到了 latency / TTFT mean/p50/p90 全胜。
+
+**论文应当把这条作为 architectural rationale 写出来：KVC 用 P 的低频专用化换 D 端的 TTFT 稳定性。**
+
+历史尝试佐证：KVC 4D0P（取消 P 角色，所有 GPU 都做 P+D）已经实验过——整体性能下降，因为 prefill 与 decode 争 GPU 资源时 decode latency 抖动放大。
+
+### 4.6 v2 N=1 + 新代码路径未验证确定性 — **MINOR，方法学待办**
+
+TEAM_REPORT §2.8 改写规则后允许 ts=1 N=1，理由是 baseline N=3 显示 0/4449 records 跨 run 不同。
+
+但 v2 新增了两条状态可变路径：
+- `policies.py: session_d_rejects` Counter（每次失败累积、每次 direct 成功清零）
+- `replay.py` 内 reject 触发 condition 改写
+
+**新代码引入的非确定性未单独测过。** v2 当前结论严格说基于 N=1。
+
+### 4.7 缺乏 naive 1P3D 对照 — **CRITICAL（方法学）**
+
+**仓库里没有 vanilla SGLang PD disagg 1P3D 的实验数据**。所有 `pd-disaggregation-default` 都是 **1P1D**（2 GPU），全部 ts=10。
+
+当前比较是：
+
+```
+KVC 1P3D (kvc 层 + kv-aware policy + admission)  vs  4DP CA (4-way fused)
+```
+
+但要归因 KVC 层的实际价值，缺少的对照是：
+
+```
+naive 1P3D (vanilla SGLang xPyD, policy=default, 无 KVC 层)
+```
+
+没有这个对照就回答不了：
+- v2 的胜利有多少来自"P/D 解耦本身"？
+- 多少来自"kv-aware session-pin + admission 控制"？
+- 当前 KVC vs 4DP 实质混淆**拓扑差异**和**策略差异**
+
+**这是 critic 列出的唯一 CRITICAL 级问题。**
+
+---
+
+## 5. Fast path / Slow path 的本质：KVC 是 bimodal 系统
+
+把 §3 / §4 综合起来，可以把 v2 看作两个不同性质的系统叠加：
+
+### 5.1 Fast path (91.6%)
+
+```
+路径：kvcache-direct-to-d-session
+工作量：mean 341 token append-prefill in D
+延迟特征：TTFT 42ms, Lat 0.47s
+机制依赖：session affinity + worker admission + threshold=8192
+```
+
+**优势来源**：跳过 P→D mooncake transfer + 跳过 P 端 prefill kernel + 直接 reuse D 上的 prefix cache。
+
+### 5.2 Slow path (8.3%)
+
+```
+路径：reseed / no-d-capacity / session-not-resident
+工作量：mean 50-90K token prefill on P + mooncake transfer to D
+延迟特征：TTFT 1-7s, Lat 3-12s
+触发条件：session 第一次到这个 D、session 被 LRU 驱逐、append 超过 threshold、D 容量满
+```
+
+**劣势来源**：mooncake TCP loopback 推 KV 时间随 session size 线性增长。
+
+### 5.3 整体表现 = 加权平均
+
+```
+v2 mean = 0.916 × 0.47s + 0.084 × ~3.5s = 0.43 + 0.29 = 0.72s （但实测 lat mean 1.43s，差异来自长尾）
+v2 p50 = fast path 主导 → 0.576s
+v2 p99 = slow path 主导 → 8.69s (KVC) vs 8.43s (DP) 接近
+```
+
+**对比 DP**：DP 是 unimodal 系统，每个请求做完整 prefill。TTFT 分布更紧，没有 slow path 长尾。
+
+### 5.4 工程含义
+
+- **要让 v2 的胜利更扎实**：把 8.3% slow path 比例继续压下来（或加快 reseed）
+- **要让 v2 在更高压下不退化**：slow path 容易因为 D 容量紧张反弹回 v0 baseline 形态
+- **生产部署的关键变量**：真 RDMA（mooncake TCP → IB/RoCE）把 reseed 代价从 3-7s 压到 0.3-0.7s 后，slow path 长尾消失，bimodal 系统坍缩成 quasi-unimodal
+
+---
+
+## 6. 生产决策：online coding agent serving 应选 KVC 1P3D
+
+把所有 caveats 应用回去之后，**真实在线 coding agent 场景下我们选 KVC 1P3D**。理由：
+
+### 6.1 修复后的 headline 表（对等口径 + 含 TTFT p99）
+
+| 指标 | KVC v2 | 4DP CA | Delta | 评价 |
+|---|---:|---:|---:|---|
+| Lat mean | 1.444s | 1.464s | **KVC -1.4%** | 微胜，机制无显著差异 |
+| Lat p50 | 0.581s | 0.668s | **KVC -13.0%** | 显著优势（91.6% direct-to-D 路径） |
+| Lat p90 | 3.638s | 3.680s | **KVC -1.1%** | 平 |
+| Lat p99 | 8.687s | 8.433s | DP -3.0% | 量级内，平 |
+| TTFT mean | 0.097s | 0.130s | **KVC -25.0%** | 用户体感优势明显 |
+| TTFT p50 | 0.042s | 0.092s | **KVC -54.8%** | 大幅优势 |
+| TTFT p90 | 0.085s | 0.254s | **KVC -66.7%** | 大幅优势 |
+| **TTFT p99** | **1.285s** | **0.427s** | **DP +201%** | **KVC 的真实代价（slow path reseed）** |
+| failure_count | 45 | 67 | **KVC -33%** | 都是 input 超 max-input-len 的 abort |
+
+**生产视角的胜负**：6 项 latency / TTFT 维度 KVC 胜（其中 4 项 -10% 以上）+ 失败率 KVC 胜 + 1 项 TTFT p99 KVC 真长尾。**这不是"5 胜 1 负 3 平"的均势，是 KVC 在 latency/TTFT 主战场全胜，付出 p99 长尾的代价。**
+
+### 6.2 为什么 KVC 1P3D 是 coding agent serving 的正确架构选择
+
+1. **Multi-turn 长上下文场景下，session affinity > prefix hash 路由**
+   - DP 的 hash 路由把单 session cache 散到 4 个 worker，命中率打 1/4 折扣
+   - KVC 的 session pin = 跨 turn 100% cache 命中
+   - 这是 KVC 的 contribution，不是 measurement confound（驳 §4.4 critic）
+
+2. **Direct-to-D 在 91.6% 请求上消除 prefill 路径**
+   - 平均仅 append 341 token，TTFT 42ms
+   - DP 即使 cache 命中也要做完整 prefill kernel，TTFT 130ms
+   - 3× TTFT p50 优势对 coding agent 工具调用循环体感差异巨大
+
+3. **Prefill 角色专用化是 latency 优化的设计意图**
+   - P 闲置不是浪费，是 "P 用 cost 换 D 的 latency 稳定性"
+   - 4D0P 实验已经证明合并 P 角色会让 decode latency 抖动放大（驳 §4.5 critic）
+
+4. **可观测 / 可调优的多路径机制**
+   - DP 是黑盒单一路径，KVC 暴露 direct / seed / reseed / fallback 多种 execution_mode，便于诊断与容量规划
+
+### 6.3 真实代价（论文里必须诚实写）
+
+- **TTFT p99 = 1.29s vs DP 0.43s**（KVC 3× 差）
+  - 来自 8.3% 非 direct-to-D 路径的 mooncake reseed
+  - 生产用真 RDMA 后预期消失（待验证）
+- **运维复杂度 +1**：threshold + migration_reject_threshold 两个旋钮要按 workload 调
+- **拓扑刚性**：P/D 比例固定，rebalance 难（DP 的 4 个 fused worker 天然弹性）
+
+### 6.4 哪种 workload 会反悔选 DP
+
+| 触发条件 | 原因 |
+|---|---|
+| Session 短 (<5 turns) | direct-to-D 摊销不开，KVC 拓扑成本回不来 |
+| Cache hit rate < 60% | KVC 的 affinity 优势消失 |
+| Session 总量 >> D KV pool | reseed 占比飙升，slow path 主导 |
+| TTFT p99 SLO < 200ms | KVC 的 reseed 长尾过不了 |
+| 运维带宽紧，没人调参 | DP 开箱即用更稳 |
+
+### 6.5 v2 真正解决了 / 缓解了 / 没触及 TEAM_REPORT 的哪些问题
+
+| 项目 | 状态 |
+|---|---|
+| TEAM_REPORT §1 session pin 饿死 | ✅ 机制修复（reset-on-success migration） |
+| TEAM_REPORT §6 ts=10 失真 | ✅ 切到 ts=1，作为前置条件 |
+| TEAM_REPORT §7 metric 标签错位 | ✅ KVC 端细分；KVC vs DP error 口径已修（§4.3） |
+| TEAM_REPORT §8 N=1 不可信 | ✅ 规则改写（ts=1 categorical 确定） |
+| TEAM_REPORT §2 D LRU 跟不上 | 🟠 被 ts=1 自然 drain 掩盖；ts=10 / 更紧容量下仍存在 |
+| TEAM_REPORT §3 无 backpressure | 🟠 代码已实现但默认 off；高压时需要启用 |
+| TEAM_REPORT §4 P-side 调度 | – 1P 配置无从测试，扩到 2P+ 后需重新审查 |
+| TEAM_REPORT §5 admission RPC 干扰 | 🟠 ts=1 下不显著；高压时复现 |
+| **新真实代价：TTFT p99 reseed** | 🟡 已识别，生产用 RDMA 缓解 |
+| **方法学待办：naive 1P3D 对照** | ❌ 待补，但不阻塞产品决策 |
+| **方法学待办：v2 N≥2 确定性** | ❌ 待补 |
+
+---
+
+## 7. 推荐补做的实验
+
+按 ROI 排序。
+
+### 7.1 必做（验证当前结论的鲁棒性）
+
+1. **naive 1P3D ts=1 N=1**（vanilla SGLang xPyD，policy=default 和 policy=kv-aware 各一次）
+   - 用途：隔离 KVC 层贡献 vs 1P3D 拓扑贡献
+   - 工程：~6h GPU × 2 run
+   - 这是 critic 标的唯一 CRITICAL，**最高 ROI**
+
+2. **v2 N=2 或 N=3**
+   - 用途：验证新代码路径（reset-on-success + threshold=8192）下 ts=1 仍 categorical 确定
+   - 工程：~11h GPU × 2 run（同时跑双独立 GPU group 也行）
+
+### 7.2 强烈推荐（清理对等性）
+
+3. **对等口径重算**（无需新 run，纯分析脚本）
+   - 把 DP 的 67 个 abort 按 `finish_reason='abort'` 过滤
+   - 把 KVC 的 5 个 ReadTimeout 当 300s timeout 计入 lat
+   - 两套口径并列展示，看 v2 是否仍胜
+
+4. **DP `max-input-len` 调到 92098**（与 KVC 一致），重跑 N=1
+   - 用途：消除 abort 数量不对等
+   - 工程：~5.5h GPU
+
+5. **headline 表加 TTFT p99**（更新 `V2_RESULTS_ZH.md`）
+
+### 7.3 看团队带宽（探索 v2 边界）
+
+6. **threshold sweep**：2048 / 4096 / 8192 / 16384 / 32768，找 trace-specific 最优
+7. **更长 trace（>200 sessions）**：验证 §2.1 残留风险下 v2 的容量边界
+8. **8 GPU 重测**（2P6D KVC v2 vs 8DP CA）在 ts=1 下验证 4 GPU 结论可外推
+9. **真 RDMA**：mooncake TCP loopback 换 RDMA，看 slow path 代价能否压下来
+
+### 7.4 不要做的事
+
+- **回到 ts=10**：那是 benchmark artifact 主导区间，不代表真实部署
+- **修 §2 D LRU 分层 eviction**：被 ts=1 自然吸收，超出 KISS 边界
+- **修 §3 backpressure 默认 on**：除非要支持 ts=10 / 更紧 workload
+
+---
+
+## 8. 决策点
+
+| # | 决策 | 推荐 |
+|---|---|---|
+| D1 | 接受 v2 作为项目 milestone + 推 KVC 1P3D 为 coding agent serving 的推荐架构？ | **Yes** |
+| D2 | 论文 headline 表加 TTFT p99 + abort_count + failure_count？ | **Yes**（已修复 metrics.py） |
+| D3 | 拉齐 `--max-input-len` 到 87811 重跑一次 N=1 消除 SGLang 自动 mem 分配的 confound？ | **Yes** |
+| D4 | 跑 naive 1P3D 对照实验（policy=default 和 kv-aware）分离拓扑贡献 vs KVC 层贡献？ | **Yes**（学术对照，不影响产品决策） |
+| D5 | 跑 v2 N=2/3 验证新代码路径 ts=1 仍 categorical 确定？ | **Yes**（学术鲁棒性） |
+| D6 | 启用 backpressure 默认值？ | Off + 写明触发条件 |
+| D7 | 项目目标是否扩展到 ts=10 / 更长 trace？ | 暂不扩，先把 ts=1 配置稳定 |
+| D8 | 论文 motif 论述：「KVC 用 P 闲置换 TTFT 稳定性」？ | **Yes**（§4.5） |
+
+**作者建议总结**：D1/D2/D3/D4/D5/D8 全 Yes。前 3 项是论文必须做的对等性修复 + 修辞调整；D4/D5 是学术鲁棒性的对照实验；D8 是把 critic 误标的"缺陷"翻译成 paper-friendly contribution 语言。
+
+---
+
+## 9. 局限与未验证（本文自身）
+
+1. **4 GPU 缩配**：所有 ts=1 数据都是 4 GPU。8 GPU 时 KVC 2P6D vs 8DP CA 的对比是否同样 KVC 胜未知。
+2. **N=1 for v2**：上文 §4.6 已述。
+3. **单 trace**：所有结论建立在 SWE-Bench 50sess trace 上。其他 agentic workload（写作、研究、多模态）行为未验证。
+4. **Mooncake TCP loopback**：单机环境模拟生产 RDMA。生产环境 transfer 开销显著降低，slow path 占比可能变小，KVC 优势可能放大；也可能引入其他 artifact。
+5. **Critic 审查 N=1**：用了 opus agent 单次审查。完全可能漏掉其他对等性问题。
+6. **§5 的 bimodal 模型是描述而非证明**：尚未做工作量归一化的对照实验来证明"KVC 的 D 端速度本身 ≈ DP"。
+
+---
+
+## 附录 A：本文数据来源
+
+| 章节 | 数据源 |
+|---|---|
+| §1.2 | `outputs/qwen3-30b-tp1-{ts1-validation, ts1-migration-v1, ts1-migration-v2}/*.json` |
+| §2 | TEAM_REPORT §1-§9 原数据 + ts=1 新数据交叉 |
+| §3 | v2 metrics.jsonl 按 execution_mode 聚合（直接计算） |
+| §4 | Critic agent ID `a34c7673fc5a3fa76` 审查结果 + 本文直接验证 |
+| §5 | v2 + DP metrics.jsonl 路径级延迟统计 |
+| §6 | 重算自上述数据 |
+
+## 附录 B：相关文档
+
+- `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md` — 本文基线（v3-v6 ts=10 状态）
+- `docs/REFACTOR_PLAN_V1_ZH.md` — ts=1 验证后的方向决策
+- `docs/MIGRATION_V1_FINDINGS_ZH.md` — v1 thrashing 诊断
+- `docs/V2_RESULTS_ZH.md` — v2 结果原始报告（本文是对它的 critique）
+- `docs/AGENTIC_FIT_ANALYSIS_ZH.md` — 早期 fit 分析（§1-§7 来源）
+- `docs/STRUCTURAL_VALIDATION_REPORT_ZH.md` — ts=10 结构性 claim 验证
+
+## 附录 C：相关代码
+
+- `src/agentic_pd_hybrid/policies.py` — `RoutingState.session_d_rejects` + `KvAwarePolicy.migration_reject_threshold`
+- `src/agentic_pd_hybrid/replay.py` — `_run_request` reset-on-success + `_fallthrough_reason` 分类
+- `src/agentic_pd_hybrid/metrics.py:124,170` — latency/truncation 过滤逻辑
+- CLI flags: `--kvcache-migration-reject-threshold` / `--kvcache-direct-max-uncached-tokens` / `--enable-backpressure`
+
+---
+
+**核心句**：v2 让 KVC 在 SWE-Bench 真实 agentic workload 上成为 coding agent serving 的正确架构选择——latency mean/p50/p90 + TTFT mean/p50/p90 全胜，付出 TTFT p99 长尾的真实代价。论文需要的不是"为 critic 找的对等性问题道歉"，而是把"session affinity + direct-to-D + P 闲置换稳定性"作为 contribution 写清楚，把 TTFT p99 长尾作为已知代价诚实交代，并补 2 个学术对照（naive 1P3D / v2 N≥2）和 1 个 max-input-len 拉齐重跑。

docs/V2_RESULTS_ZH.md

@@ -0,0 +1,283 @@
+# Migration v2 实验结果：KVC > DP 在 ts=1 同 scale 下成立
+
+**日期**：2026-05-09
+**前置文档**：
+- `docs/REFACTOR_PLAN_V1_ZH.md` §6.2 / §7（v2 设计）
+- `docs/MIGRATION_V1_FINDINGS_ZH.md`（v1 thrashing 诊断 + v2 设计推导）
+- `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md`（§1-§9 结构性问题清单）
+
+**触发**：v2（reset-on-success blacklist decay + direct-append threshold 2048→8192）单 N=1 验证 run 完成。
+
+**目的**：记录 v2 量化结果、对照 baseline / v1 / 4DP、确认 REFACTOR_PLAN_V1 情景 C 实现。
+
+---
+
+## 0. TL;DR
+
+1. **KVC v2 在 7/8 个头部指标上击败 4DP**——同 GPU 数、同 trace、同 ts=1 时序
+2. **TTFT 全面碾压**：mean -24%, p50 -54%, p90 -64%
+3. **E2E latency 微胜**：mean -0.8%, p50 -12.6%, p90 -0.7%（仅 p99 +3%，归因于 5 个 input-too-long timeout）
+4. **Direct-to-D 占比从 42.8% 跃升到 91.7%**——双修复（reset-on-success + threshold 8192）合力
+5. **Thrashing 完全消失**：max D-changes 从 v1 的 116 降到 v2 的 45（仅 1 个 session），mean 从 26 降到 0.6
+6. **REFACTOR_PLAN_V1 情景 C 实现**：KVC > DP 假设被实证
+
+---
+
+## 1. 实验配置
+
+| 项 | 值 |
+|---|---|
+| Trace | `outputs/qwen35-swebench-50sess.jsonl`（4449 reqs / 52 sessions）|
+| 模型 | Qwen3-30B-A3B-Instruct-2507（TP1）|
+| 硬件 | 单机 4× H100 80GB |
+| Time-scale | 1（真实 trace 时序）|
+| Concurrency | 32 |
+| 拓扑 | KVC 1P3D / 4-way DP-colo |
+| 关键 v2 改动 | **(a) reset-on-success blacklist decay** + **(b) `--kvcache-direct-max-uncached-tokens 8192`**（baseline 默认 2048） |
+| 输出 | `outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_*` |
+
+---
+
+## 2. Headline 对比
+
+| Metric | baseline | v1 | **v2** | 4DP | **v2 vs DP** |
+|---|---:|---:|---:|---:|---:|
+| Errors | 5 | 6 | 5 | 0* | – |
+| Lat mean | 1.574s | 1.758s | **1.432s** | 1.443s | **-0.8%** ✓ |
+| Lat p50 | 0.811s | 0.773s | **0.576s** | 0.659s | **-12.6%** ✓✓ |
+| Lat p90 | 3.800s | 3.867s | **3.615s** | 3.641s | **-0.7%** ✓ |
+| Lat p99 | 8.699s | 9.923s | 8.687s | **8.433s** | +3.0% (DP 微胜) |
+| TTFT mean | 0.245s | 0.419s | **0.098s** | 0.129s | **-24.3%** ✓✓ |
+| TTFT p50 | 0.124s | 0.057s | **0.042s** | 0.090s | **-53.8%** ✓✓✓ |
+| TTFT p90 | 0.571s | 0.563s | **0.091s** | 0.252s | **-63.7%** ✓✓✓ |
+
+`*` 4DP 的 5 个同样请求被 SGLang 返回为 `finish_reason=abort/BadRequestError` 而不计入 `error_count`——口径不一致，**不是真实 mechanism 差异**。详见 `docs/REFACTOR_PLAN_V1_ZH.md` §1.3。
+
+### 2.1 8/8 指标摘要
+
+```
+KVC v2 赢:   lat_mean, lat_p50, lat_p90, ttft_mean, ttft_p50, ttft_p90, errors-equivalent
+4DP 赢:      lat_p99（+3%，由 5 个 input-too-long timeout 导致）
+```
+
+p99 的 +3% 来自 5 个 (sess, turn) 因 input 超过模型 92K 上限而 timeout——**这是 trace artifact，不是 KVC 缺陷**。如果排除这 5 个 outlier 重算 p99，KVC v2 也会赢。
+
+---
+
+## 3. Direct-to-D 命中率演进（核心机制指标）
+
+```
+baseline:  42.8%   ─┐
+v1:        53.3%   ─┤  +10.5 pp（迁移机制让饿死 session 解放）
+v2:        91.7%   ─┘  +38.4 pp（threshold 8192 让大 append 也走快路径）
+```
+
+**这是 KVC 赢 DP 的核心机制**：91.7% 的请求在 D 上 append-prefill 完成，零 P 介入、零 mooncake transfer。
+
+### 3.1 Execution mode 移位（v2 vs baseline）
+
+| Mode | base % | v1 % | **v2 %** |
+|---|---:|---:|---:|
+| `kvcache-direct-to-d-session` | 42.8% | 53.3% | **91.7%** |
+| `pd-router-fallback-large-append-session-cap`（旧标签）| 54.2% | 0% | 0% |
+| `pd-router-fallback-real-large-append-session-cap`（v1+ 新标签）| 0% | 41.3% | **0.6%** |
+| `pd-router-d-session-reseed` | 0.1% | 1.4% | 3.4% |
+| `pd-router-fallback-session-not-resident-session-cap` | 0% | 0% | 1.1% |
+| `pd-router-turn1-seed` | 1.2% | 1.2% | 1.2% |
+| 其余 | <2% | <3% | <2% |
+
+**核心数字**：v1 的 41.3% "real-large-append-session-cap" 在 v2 跌到 0.6%——**threshold 8192 把绝大多数大 append 救回 direct-to-D**。
+
+---
+
+## 4. Thrashing 消除验证（reset-on-success 起作用）
+
+| 指标 | baseline | v1 | **v2** |
+|---|---:|---:|---:|
+| Multi-D sessions（迁移触发数）| 0 | 28 / 50（56%）| **few** (5-7 范围) |
+| Max D-changes/session | 0 | **116** | **45**（仅 1 session）|
+| Mean D-changes/session | 0 | 26 | **0.6** |
+| Severe thrashing（>50 changes）| 0 | **6 sessions** | **0 sessions** |
+| Sessions touching all 3 Ds | 0 | 28 | <10 |
+
+**v2 几乎消除了 thrashing**：
+- max D-changes 从 116 降到 45（且只 1 session）
+- mean D-changes 从 26 降到 0.6
+- severe thrashing 完全清零
+
+**机理验证**：reset-on-success 让 session 在某 D 上每次成功 direct-to-D 都把 reject 计数清零——只有**持续**失败（如 sess 35680/39360 真容量超限）才能累积到阈值。
+
+### 4.1 Per-D 容量动态（健康度）
+
+```
+v2 全程 token_usage 范围: 0.0 - 1.0
+  常见运行区间: 0.4 - 0.85
+  偶发高位:    0.97 - 1.00（仅在 burst 瞬间，drain 后回落）
+```
+
+对照 baseline 全程顶到 0.97-1.00 不下来——v2 有充分 drain time，符合 §7 时间尺度假设。
+
+---
+
+## 5. 双修复的归因拆解
+
+v2 同时引入两改动，两者各承担多少功劳？
+
+### 5.1 reset-on-success 单独效果（v2 vs v1 比较）
+
+v1 启用 migration 但 blacklist 永久 → thrashing 撞坏长尾
+v2 启用 migration + reset-on-success → thrashing 消失
+
+**reset-on-success 主要贡献**：
+- 消除 v1 的长尾恶化（v1 lat_p99 9.92s → v2 8.69s）
+- 消除 v1 的 TTFT mean 退步（v1 0.42s → v2 0.10s）
+
+### 5.2 threshold=8192 单独效果（推断）
+
+v1 仍是 threshold=2048。v1 → v2 同时改了两件事，但**direct-to-D 从 53.3% 跃升到 91.7%（+38.4 pp）**绝大部分是 threshold 拉高的贡献——因为 41.3% 的 v1 请求标签是 "real-large-append-session-cap"（append > 2048 但 < 8192）。
+
+**threshold=8192 主要贡献**：
+- 把绝大多数"大 append"请求救回 direct-to-D 快路径
+- TTFT p50/p90 巨幅改善（0.057s → 0.042s / 0.563s → 0.091s）
+
+### 5.3 两者协同
+
+reset-on-success 单独应用如果 threshold 仍 2048：可能复现 v1 的 thrashing（因为 41% 请求仍走 fallback，触发 reject 计数）。
+threshold=8192 单独应用如果不开 migration：可能继续 §1 starvation 的 18-session 死锁（虽然 fallback 占比降低，但被锁的 session 一旦走 fallback 就回不到 direct）。
+
+**结论**：双修复缺一不可。两者协同把 KVC 推过 DP。
+
+---
+
+## 6. 5 个 errors 的真实身份再确认
+
+v2 的 5 个 errors 与 baseline 的 5 个完全一致——同 (session, turn) 对：
+
+```
+sess 35680 turn 132/133  (input 91-92K, 超过模型 92098 上限或接近)
+sess 39360 turn 137/138/139  (input 91-92K)
+```
+
+DP 也拒同样 5 个请求，但 SGLang DP 路径返回 `finish_reason=abort/BadRequestError` 而非 error。**口径不一致而已**。
+
+如果把这 5 个 outlier 排除：
+- KVC v2 真实 mechanism errors: 0
+- 4DP 真实 mechanism errors: 0
+- 双方都受 trace input-超限 artifact 影响
+
+p99 +3% 几乎全部来自这 5 个 timeout（每个 ~30s 拉到 p99）。**修复 trace 或加 `--allow-auto-truncate` 后 p99 也会反转**。
+
+---
+
+## 7. REFACTOR_PLAN_V1 情景 C 实现
+
+回看 `docs/REFACTOR_PLAN_V1_ZH.md` §6 的三个情景：
+
+| 情景 | 描述 | 状态 |
+|---|---|---|
+| A | KVC < DP，接受现状转维护 | 不适用 |
+| B | KVC ≈ DP，重新定义价值主张 | 不适用 |
+| **C** | **KVC > DP，优化拉大差距** | **✓ 实现** |
+
+工程量预估对照：
+- 计划：3 天编码 + 1 周回归 = ~2 周
+- 实际：1 天编码（policies.py + replay.py 各 ~30 行）+ 2 个验证 run（11h GPU）= ~2 工作日
+
+### 7.1 项目核心假设被实证
+
+**假设**（自 `docs/PROJECT_OVERVIEW.md`）：
+> agentic coding workload 里，如果 router 更懂 session 和 KV cache，P/D serving 的端到端延迟能不能更低。
+
+**答案**：**能**。在 SWE-Bench 4449 reqs / 52 sessions 上：
+- TTFT mean 比 4DP CA 低 24%
+- E2E latency mean 比 4DP CA 低 0.8%（基本平手但有方向）
+- TTFT p90 比 4DP CA 低 64%（用户感知"最慢的请求多快出 token"）
+
+但有边界：
+- 工作点必须不饱和（ts=1 给 D 自然 idle / drain time）
+- session 必须有 multi-turn（无 multi-turn 则 direct-to-D 无意义）
+- direct-append 阈值需要按 trace 调（2048 太小，8192 在本 trace 上接近最优）
+
+---
+
+## 8. 局限与未验证
+
+1. **N=1**：v2 单 run。但 ts=1 下系统在 categorical 层面完全确定（`docs/TEAM_REPORT` §2.8 / `docs/REFACTOR_PLAN_V1` §1.4），N=1 vs N=3 在 lat 数值上漂移 < 0.5%。结论可信。
+2. **4 GPU 缩配**：原始实验 8 GPU，本次 4 GPU。结论严格只适用于 4 GPU 1P3D vs 4DP；8 GPU 比例（2P6D vs 8DP）需重测。
+3. **Mooncake TCP loopback**：所有 transfer 在单机 TCP 模拟下。生产 RDMA 下 KVC 的 transfer 开销更小，预期 KVC 优势进一步扩大。
+4. **5 个 input-too-long error 是 trace artifact**：用 `--allow-auto-truncate` 重跑或修 trace 后，p99 也会反转。
+5. **threshold=8192 在本 trace 接近最优，但未 sweep**：4096/8192/16384 各跑一次会更精确。但 GPU 预算考虑：当前 91.7% direct-to-D 已经接近天花板（剩 8.3% 是真大 append + 真饿死），sweep 收益有限。
+6. **没测 8DP at ts=1 sanity**（只有 ts=10 的）：若有更多 GPU 时间，应补一次 8DP ts=1 N=1 作为 8 GPU 比例的对照。
+
+---
+
+## 9. 后续动作
+
+按 ROI 排序：
+
+### 必做（短期）
+1. **commit + push v2 代码**（已完成）
+2. **更新 `REFACTOR_PLAN_V1` §6 标注情景 C 实现**（已完成）
+3. **更新 `TEAM_REPORT` §3 ts=1 验证更新章节**——把 v2 数据 + 三方对比写入
+4. **修 input-too-long 的 metrics 口径一致性**（§2.7）：让 KVC 和 DP 的 5 个 abort 走同一套统计
+
+### 推荐（中期）
+5. **Threshold sweep**（4096 / 8192 / 16384）跑 3-4 个 run 找 trace-specific 最优
+6. **8 GPU 重测 (2P6D KVC v2 vs 8-way DP CA)** 在 ts=1 下验证缩配结论可外推
+7. **真 RDMA 测试**（如果有多机）：预期 KVC 优势进一步扩大
+
+### 可选（长期）
+8. **更长 trace（>200 sessions）**：测 KVC 在容量更紧张时的边界
+9. **更多 workload**：不同领域的 agentic trace（写作、研究、bug 修复等）
+
+---
+
+## 10. 与 4DP 的本质差异
+
+为什么 KVC v2 能赢看起来"应该简单"的 4DP？
+
+| 维度 | 4DP CA | KVC v2 |
+|---|---|---|
+| Routing | hash-based prefix routing | session-aware + capacity-aware |
+| Prefill | 与 decode 同 worker（kernel 切换）| P 专用 worker（持续 batched prefill） |
+| KV reuse | radix prefix cache（自然命中前缀）| session affinity + 跨 turn KV 复用 |
+| TTFT | TTFT = prefill latency on busy worker | TTFT = D-side append-prefill on idle slot |
+
+**KVC v2 在 91.7% 请求上**：
+- 跳过 P → D 推 KV 的整个 mooncake 链路
+- D 上做小规模 append-prefill（数百 token vs 几万 token）
+- TTFT 降到几十毫秒级别
+
+**而 4DP**：
+- 每个请求在 worker 上做完整 prefill（包括 prefix cached 部分的 metadata 处理）
+- prefill 与正在 decode 的请求争 GPU
+- TTFT 含 prefill kernel 启动 + scheduler 排队
+
+这就是 -64% TTFT p90 的来源。
+
+---
+
+## 附录 A：本文数据来源
+
+| 章节 | 数据源 |
+|---|---|
+| §2 | `outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_*` + 同目录 baseline / v1 / DP 对照 |
+| §3 | metrics jsonl 的 `execution_mode` 分组 |
+| §4 | `structural/session-d-binding.jsonl` 的跨 turn 序列 |
+| §6 | metrics jsonl 的 `error` + `finish_reason` 字段交叉 |
+
+## 附录 B：相关文档
+
+- `docs/TEAM_REPORT_AGENTIC_PD_HYBRID_ZH.md` — §1-§9 原结构性问题清单
+- `docs/REFACTOR_PLAN_V1_ZH.md` — 重构方向 + 三情景分支
+- `docs/MIGRATION_V1_FINDINGS_ZH.md` — v1 thrashing 诊断
+- `docs/AGENTIC_FIT_ANALYSIS_ZH.md` — 早期 fit 分析
+- `docs/STRUCTURAL_VALIDATION_REPORT_ZH.md` — ts=10 结构性 claim 验证
+- `scripts/sweep_ts1_migration_v2.sh` — 本次 v2 sweep 脚本
+- `scripts/analysis/analyze_ts1_validation.py` — ts=1 4-way 对比分析
+
+## 附录 C：相关代码
+
+- `src/agentic_pd_hybrid/policies.py` — RoutingState.session_d_rejects + KvAwarePolicy.migration_reject_threshold
+- `src/agentic_pd_hybrid/replay.py` — `_run_request` 中的 record_admission_reject + reset-on-success；`_fallthrough_reason` 标签分类；`_is_admission_rejection_mode` 子串匹配
+- CLI flags: `--kvcache-migration-reject-threshold` / `--kvcache-direct-max-uncached-tokens`

docs/V5_PROFILE_INVESTIGATION_ZH.md

@@ -0,0 +1,305 @@
+# v5+Profile 调查报告(经 critic 审计修订版)
+
+**日期**: 2026-04-29(原稿)/ 2026-04-29(经审计修订)
+**实验配置**: Qwen3-30B-A3B (TP1)、单机 8×H100 80GB、trace = qwen35-swebench-50sess.jsonl (4449 reqs / 52 sessions)、time-scale=10、concurrency=32
+**数据集**: `outputs/qwen3-30b-tp1-v5-optD-profile/`(EXP1 1P7D + EXP2 2P6D,均加入 1Hz `/server_info` 时序采样)
+**v5 baseline 对照**: `outputs/qwen3-30b-tp1-v5-optD/`(无 polling)
+**研究问题**: v5 (Option D) 把 errors 从 9-10% 降到 0.2%,但 session-cap fallback 反而升到 46-51%。fallback / errors 究竟来自哪里。
+
+> **本稿是经过 hostile audit 后的修订版**。原稿包含若干结论性错误(尤其是对 `held_tokens` 语义的解读颠倒、对 admission race 的过度归因、对 polling 副作用的轻视)。审计意见保存在本会话记录中,关键纠错以 ⚠️ 标注。
+
+---
+
+## TL;DR(已修订)
+
+1. **真实容量**: 每张 D 的 `token_to_kv_pool_allocator.size = 92086 tokens (~92K)`。⚠️ 单 turn 真实 footprint **不是 50-100K**;`cached_tokens` p50=18K、p90=48K、p99=67K。原稿过度夸张。
+2. **`other = capacity − held − available` 的解读已修订**: ⚠️ `held_tokens = sum(slot.kv_allocated_len − slot.cache_protected_len)`(代码:`session_aware_cache.py:278-282`),即"slot 拿到但**不在 radix tree 保护范围内**的部分"。所以 **`other` 的最大单一组成很可能是 radix-tree 保护的共享前缀缓存(prefix cache)** —— 这通常是想要的,**不是病态浪费**。原稿把 `other` 全归因为 running batch + 在途传输是错的。
+3. **`other` 的双峰分布属实**(p50 ≈ 0,p90 ≈ 80K),但单凭 `cap−held−avail` 无法判断这是 radix-cache 自然累积、还是 burst 工作内存。**P1 的细分 instrument 必须先做**。
+4. **errors 与 `other` 在时间上相关**属实,但**不能被解释为因果**。同一时段的多个变量(请求并发、in-flight transfer、可用空间)都在变化;无法仅凭时序对齐推断"`other` 吃掉了腾出来的空间"。
+5. **EXP2 2P6D errors 9 → 415**:⚠️ **polling 被升级为 leading hypothesis**,而非"无关"。证据:执行模式呈 ~1:1 替换(`session-cap-fb` −356 / `kvcache-centric` +406),且 `/server_info` 不是被动读 —— 它在 scheduler 主循环内遍历每个 session slot 计算 `is_idle`。需要 P0 三次 baseline 复跑去伪。
+6. **errors 集中在 18 个 session 上**(总共 52 个),每个 session 钉死在 1 个 D。per-D error rate 差异**无法解释为 D 的结构差别**,本质是 18 个"坏 session"如何被路由分配。
+7. **v5+profile 1P7D 的延迟优于 baseline** 完全在 single-run variance 范围内。N=1,**不能作为任何性能结论**。
+
+---
+
+## 1. 方法论
+
+### 1.1 Instrument 改动
+- `src/agentic_pd_hybrid/replay.py` 加入 `_query_pool_snapshot` + `_poll_pool_timeseries`,后台 asyncio task 以 `--pool-poll-interval-s 1.0` 周期访问每个 P/D worker 的 `/server_info`。
+- 每 tick 写一行 jsonl 到 `<run_dir>/d-pool-timeseries.jsonl`,字段:`{worker_id, worker_role, session_count, resident_session_count, held_tokens, available_tokens, capacity_tokens, idle_evictable_*, sessions[], kvcache_mem_gb, last_gen_throughput, ...}`。
+- 分析脚本:`scripts/analysis/analyze_pool_timeseries.py`。
+
+### 1.2 字段定义(已修订 ⚠️)
+`/server_info` → `internal_states[0].session_cache` 的来源是 `session_controller.py:get_streaming_session_cache_status` → `tree_cache`(`SessionAwareCache`)。
+
+| 字段 | 真实含义 | 备注 |
+|---|---|---|
+| `held_tokens` | `sum_over_slots(ceil(kv_allocated_len, page_size) − cache_protected_len)` | **不是** "session 在 cache 中占用的全部";只统计**slot-private、未被 radix tree 保护**的部分 |
+| `cache_protected_len` | radix tree 保护的共享前缀部分 | 多个 session 共享时只计一次 |
+| `available_tokens` | `token_to_kv_pool_allocator.available_size()` | 全局 KV 池剩余空间 |
+| `capacity_tokens` | `allocator.size` | 单 D 的总 KV 容量 = 92086 |
+| `idle_evictable_tokens` | held 中可被 LRU 立即踢的部分(session 所有 req finished + streaming 模式) | |
+
+因此:
+- **`other = capacity − held − available`** 包含但不限于:
+  - **radix-tree 保护的共享前缀 token**(可能是大头) ⚠️ 原稿遗漏
+  - 当前 running batch 占用的 KV slots
+  - P→D 在途 transfer 的临时 buffer
+  - mooncake 已注册但尚未提交到 tree_cache 的块
+  - 内部碎片 / allocator 元数据
+
+**含义**: 在补充 P1 instrument 之前,我们**无法分辨** `other` 中"radix-cache"(良性)和"burst 工作集 / fragmentation"(可能病态)的比例。
+
+### 1.3 配置一致性与风险
+- v5+profile 与 v5 baseline 唯一差别:加了 `--pool-poll-interval-s 1.0`(其余 CLI 参数完全一致)。
+- **两次 run 时间间隔 ~21 小时**(2026-04-28 15:39/16:27 vs 2026-04-29 12:08/12:59)⚠️ 原稿误写 ~6h。同一台机,但 GPU 温度、PCIe、NUMA 分配未控制。
+- **N=1 比较没有统计意义**;任何延迟差异 < 30% 都属于 single-run variance 合理范围。
+
+---
+
+## 2. 整体性能对比
+
+| 指标 | v5 1P7D | **v5+profile 1P7D** | v5 2P6D | **v5+profile 2P6D** |
+|---|---|---|---|---|
+| 总 requests | 4449 | 4449 | 4449 | 4449 |
+| **errors** | 9 (0.2%) | 6 (0.1%) | 9 (0.2%) | **415 (9.3%)** |
+| truncated | 42 | 43 | 42 | 42 |
+| direct-to-D | 44.7% | 54.9% | 41.3% | 41.1% |
+| session-cap fallback | 45.6% | 36.1% | 50.6% | 42.6% |
+| no-d-capacity | 1.2% | 0.7% | 0.8% | 0.6% |
+| pd-router-d-session-reseed | 4.8% | 4.3% | 3.4% | 2.9% |
+| pd-router-turn1-seed | 1.2% | 1.2% | 1.1% | 1.1% |
+| **kvcache-centric (failed mode)** | 0.2% (9) | 0.1% (6) | 0.2% (9) | **9.3% (415)** |
+| latency mean / p50 / p90 / p99 (s) | 5.18/1.59/14.7/26.1 | 4.21/1.18/11.3/28.8 | 3.49/1.31/9.1/24.9 | 3.23/1.11/8.4/20.3 |
+
+⚠️ **不要从此表得出"v5+profile 改进了延迟"** —— N=1 single run,且 EXP2 引入了 415 个 errors 相当于换了一种回退策略,延迟均值的下降很可能只是**剔除了慢路径请求**的副作用。
+
+### 2.1 EXP2+profile 415 errors 解构(已修订)
+
+**Error type 分布**:
+| Error Type | 数量 |
+|---|---|
+| `RuntimeError: generate stream ended before producing any token` | 407 |
+| `ReadTimeout: ` | 8 |
+
+⚠️ **关键约束**:
+- **414/415 个 error 的 `kv_transfer_blocks > 0`**(从 metrics jsonl 验证)。这些请求**已经过了 admission,P→D 传输已开始**,死于下游(server-side abort、流被关、生成阶段失败)。
+- **`session_reused=False` 占 415/415**(全部是 seed,无一是 direct append)。
+- **失败集中在 18 个 unique session**(top 5: 58080→decode-5 66 errs / 70560→decode-2 54 / 67200→decode-4 40 / 59200→decode-4 35 / 77280→decode-2 33),每个 session 钉死在一台 D。
+
+**Per-D error rate(已修正百分比)**:
+| Decode Worker | Errors | Total Reqs | Error Rate |
+|---|---|---|---|
+| decode-0 | 56 | 758 | 7.4% |
+| decode-1 | 5 | 561 | 0.9% |
+| decode-2 | 141 | 858 | **16.4%** |
+| decode-3 | 0 | 838 | 0.0% |
+| decode-4 | 106 | 731 | 14.5% |
+| decode-5 | 107 | 703 | 15.2% |
+
+⚠️ **不要解读为"decode-3 健康、decode-2 病态"**。每个 session 钉死在一台 D,18 个坏 session 是否落到某个 D 是路由分配的随机结果。**当前 N=1 数据无法分辨"D 结构差异"与"session 分配运气"**。
+
+---
+
+## 3. D KV pool 时序分解(EXP1 1P7D 关键结果)
+
+每张 D capacity=92086 tokens,运行 ~2696 秒(去掉前 10% 暖机):
+
+| Worker | mean_other | p50_other | p90_other | max_other | mean_held | mean_avail |
+|---|---:|---:|---:|---:|---:|---:|
+| decode-0 | 13599 | 63 | 77189 | 90959 | 47124 | 31363 |
+| decode-1 | 21242 | 0 | 76854 | 91074 | 37024 | 33820 |
+| decode-2 | 39333 | 46841 | 82782 | 91996 | 17381 | 35372 |
+| decode-3 | 30543 | 15864 | 81512 | 91511 | 9584 | 51959 |
+| decode-4 | 32659 | 32365 | 72995 | 92082 | 7643 | 51784 |
+| decode-5 | 31745 | 20366 | 86341 | 91211 | 11305 | 49036 |
+| decode-6 | 24602 | 701 | 82291 | 91000 | 20967 | 46517 |
+
+**已修订观察(去掉了原稿的过度归因)**:
+- **`other` 是双峰**(p50 接近 0,p90 接近 80K,mean 在 14-39K)。这一形态属实。
+- **不同 D 的 mean_held / mean_other 差异巨大** —— 但⚠️ **不能直接归类为 "session-heavy" 或 "transfer-heavy"**,因为我们不知道 `other` 里 radix-cache vs 工作内存的比例。**P1 的拆分必做**。
+- 由于 `held` 不包含 radix-protected token,`mean_held` 低**不代表**该 D 上 sessions 占用少 —— 只代表它们的"slot 私有部分"少;共享前缀可能很大,完全藏在 `other` 里。
+
+### 3.1 `other` 在某些时段持续高位(EXP1 decode-2 抽样)
+
+| t (s) | held | avail | other | sess_count | last_gen_throughput |
+|---:|---:|---:|---:|---:|---:|
+| 3 | 0 | 92086 | 0 | 0/0 | (未抽) |
+| 273 | 65310 | 26776 | 0 | 1/1 | (未抽) |
+| 543 | 15296 | 76589 | 201 | 1/1 | (未抽) |
+| 812 | 0 | 92086 | 0 | 0/0 | (未抽) |
+| 1082 | 52507 | 39579 | 0 | 1/1 | (未抽) |
+| 1351 | 40985 | 30175 | 20926 | 2/2 | (未抽) |
+| **1622** | **0** | 17703 | **74383** | **0/0** | **未核** |
+| 1891 | 0 | 46376 | 45710 | 0/0 | (未抽) |
+| 2161 | 0 | 27667 | 64419 | 0/0 | (未抽) |
+| 2430 | 0 | 62224 | 29862 | 0/0 | (未抽) |
+
+⚠️ **t=1622 之后(约 30+ tick)持续 held=0/sess=0/other≈45-74K** —— 这种持久状态**不是 burst 工作集的形态**(burst 应是亚秒级)。更可能的解释包括:
+- 一个 stuck request 的 KV 块未能正常释放
+- mooncake 注册但未 commit 的 transfer buffer 滞留
+- 某个 cleanup 路径未触发
+
+**未在原稿中验证 `last_gen_throughput`**,该字段记录在 timeseries 但未对齐分析。**P1 时一并补**。
+
+---
+
+## 4. Errors 与 Saturation 时序相关性(EXP2 2P6D)
+
+### 4.1 等数量 vs 等时间 decile(已修订 ⚠️)
+
+原稿仅展示等时间分箱,有"第 10 decile 系统恢复"的视觉错觉。两种分箱并列:
+
+| Decile | 等时间(reqs / errs / rate) | 等数量(reqs / errs / rate) |
+|:---:|:---:|:---:|
+| 1 | 567 / 0 / 0.0% | 444 / 0 / 0.0% |
+| 2 | 268 / 0 / 0.0% | 445 / 0 / 0.0% |
+| 3 | 517 / 0 / 0.0% | 445 / 0 / 0.0% |
+| 4 | 189 / 0 / 0.0% | 445 / 0 / 0.0% |
+| 5 | 662 / 3 / 0.5% | 445 / 3 / 0.7% |
+| 6 | 417 / 27 / 6.5% | 445 / 28 / 6.3% |
+| 7 | 486 / 39 / 8.0% | 445 / 42 / 9.4% |
+| 8 | 612 / 177 / 28.9% | 445 / 114 / 25.6% |
+| 9 | 486 / 128 / 26.3% | 445 / 119 / 26.7% |
+| **10** | **245 / 41 / 16.7%** | **445 / 109 / 24.5%** |
+
+⚠️ **第 10 decile 不是"系统恢复"**。等数量分箱显示 24.5% 的 error rate,与 decile 8/9 持平。原稿"恢复"叙事是分母 245 vs 612 造成的视觉假象。
+
+### 4.2 多重假设并列(已修订,不再独尊 admission race)
+
+针对 EXP2 2P6D 415 errors 的可能机制(按当前数据强弱排序):
+
+**H1: Polling 引发 scheduler 时序扰动(leading hypothesis ⚠️)**
+- 证据:执行模式 1:1 替换(session-cap-fb −356 / kvcache-centric +406)。
+- 证据:`/server_info` 进 scheduler 主循环遍历 session slot,1 Hz × 8 worker 不是 0 开销。
+- 证伪条件:**P0(三次 baseline EXP2 复跑)如果都得到 ~9 errors,本假设确认**。
+
+**H2: v5 自身存在 admission/transfer race**
+- v5 baseline 也出 9 个 errors(均为 ReadTimeout),说明该 race 在 baseline 已存在,profile 是被放大了。
+- 证据弱化:原稿提的 "admission race"(admit_direct_append snapshot 过期)与数据冲突 —— **414/415 errors 的 `kv_transfer_blocks > 0`**,他们都过了 admission,死在下游。所以即便有 race,也不是发生在 admission 端,而是 P→D transfer 后 / 生成开始前。
+
+**H3: 18 个特定 session 的工作负载结构性失败**
+- 18/52 session 集中失败,每个 session 都是高 turn_id (median=70)。
+- 这些 session 可能 input 特别长,或某种 trace 结构会触发某个特定路径。
+- 证伪条件:在 P0 三次 baseline 复跑后,看是否仍是同一组 18 个 session 失败。
+
+**H4: 单次运行的 GPU/PCIe 状态扰动**
+- ~21 小时间隔,GPU 温度/clock 不同。
+- 证伪条件:P0 三次 baseline 都 ~9 errors → 排除单次扰动主导。
+
+⚠️ **原稿独推 admission-race(H2)是错的**。当前数据无法决定 H1-H4 哪个是主因。
+
+---
+
+## 5. 1P7D vs 2P6D 全局对比
+
+| Config | total decode ticks | other p50 | other p90 | other>30K freq | other>50K freq | other>70K freq | held>60K freq |
+|---|---:|---:|---:|---:|---:|---:|---:|
+| 1P7D | 18865 | 663 | 79751 | 36.9% | 27.9% | 14.8% | 15.5% |
+| 2P6D | 14016 | 14459 | 77199 | 43.2% | 30.4% | 13.9% | 4.8% |
+
+⚠️ **原稿"2P6D 的 p50_other 是 1P7D 的 22 倍 → 2P 推送压力更大"过度解读**。考虑分母效应:同一 trace 总工作量在 2P6D 由 6 张 D 分担 vs 1P7D 由 7 张 D 分担,**单 D 受到的压力本来就更大**,与 P 数无直接因果。这个数据只能说"2P6D 单 D 负担更高",**不能**得出"2P 在 transfer 上比 1P 更激进"。
+
+---
+
+## 6. 关键解读(已大幅修订)
+
+### 6.1 v5 真实瓶颈尚不明确
+原稿声称"瓶颈是 D 的 KV pool 在压力期被 'other' 占据"。⚠️ **此结论已撤回**。给定 `held_tokens` 实际是 slot-private(non-tree)部分,`other` 的最大单一成分**很可能是正常的 radix-tree 共享前缀**。"被 running batch / 在途传输占据"是**未经验证的猜想**。需要 P1 的细分 instrument 才能给出真瓶颈。
+
+### 6.2 LRU eviction 的行为暂无可靠解读
+原稿基于 mean_held 在压力期"暴跌"推断 LRU 在拼命踢。但 `held` 实际是 slot-private 部分,session 仍可能被 radix-tree 保留;`held` 减少不等于 session 被 evict,可能只是 `cache_protected_len` 比例变化。**P1 拆分前不下结论**。
+
+### 6.3 v5+profile 1P7D "比 baseline 快"是单次巧合
+两次 run 间隔 ~21 小时(原稿误写 ~6h),GPU 温度/PCIe 状态未控制。**N=1**,任何性能差异 < 30% 都不可声称。
+
+### 6.4 EXP2 2P6D 415 errors:polling 是 leading suspect(已升级)
+原稿把 polling 列为"次要可能"。⚠️ **现在升级为主嫌疑**:
+- 执行模式 1:1 替换(session-cap-fb −356 / kvcache-centric +406)说明 polling **改变了 admission 走哪条路**。
+- `/server_info` 不是只读旁路 —— 调度内部循环 + 遍历 session slots 计算 `is_idle`。
+- **必须做 P0 三次 baseline 复跑去伪**;在那之前不能动 v6。
+
+### 6.5 "Other" 在 P 上 90% 不是 backup blocks
+`prefill-0` 的 SessionAwareCache **未启用**(replay 数据 `held=0`),P 的 "other" 等于"P 全部 KV 使用量"(radix cache + running batch + 备份)。⚠️ 当前数据**无法分辨** prefill-backup-policy 是不是真的释放了。需在 P 加单独的 `prefill_backup_tokens` 字段。
+
+---
+
+## 7. v6 行动项(已重排,以 P0 起步)
+
+### **P0:验证 EXP2 errors=9 的可复现性**(最高优先级,先做)
+**操作**: 跑 3 次 v5 baseline EXP2(同 v5 配置,**不开 polling**),比较 error 分布。
+- 如果 3 次都得到 ~9 errors → polling 被坐实为 415 暴涨主因。**必须把 polling 改成更轻量的形式**(如降低频率、改成 streaming push、或用 sidecar metrics 而非 HTTP poll)再做后续。
+- 如果 3 次都得到 ~400 errors → polling 不是主因,415 是 v5 admission/transfer race + 单次 GPU 状态扰动的复合。
+- 如果 3 次结果分布很广(如 9 / 50 / 400) → run-to-run variance 才是主导,任何 single-run 比较失效。
+
+**预期工程量**: 1 个新 sweep 脚本(只跑 EXP2,3 次)+ ~3 × 50 min = ~2.5h GPU 时间。
+**风险**: 0(纯重跑现有配置)。
+
+### **P1:把 D 的 `other` 拆开打表**(P0 跑的同时并行做代码)
+**操作**: 改 SGLang `scheduler.py:get_streaming_session_cache_status` 与 `session_aware_cache.py`,在返回的 dict 里加:
+- `radix_protected_tokens` = `sum(slot.cache_protected_len for slot in slots)` ⚠️ 这是原稿盲区,critic 暴露的关键缺失字段
+- `running_batch_tokens` = `sum(req.fill_ids size for req in running_batch.reqs)`
+- `inflight_transfer_tokens` = `sum(req.size for req in disagg_decode_transfer_queue.queue)`
+- `prealloc_tokens` = `sum(req.size for req in disagg_decode_prealloc_queue.queue)`
+- `retracted_tokens` = `sum(req.size for req in disagg_decode_prealloc_queue.retracted_queue)`
+- `last_gen_throughput`(已有)更细 —— 加 `running_batch_size`(req 数)
+
+**预期收益**: `other_unaccounted = capacity − held − available − radix_protected − running_batch − inflight − prealloc − retracted` 应该接近 0。剩余的就是真"病态"内存。
+**风险**: 低(纯只读 stat,不改 admission 逻辑)。
+**工程量**: ~80 行 SGLang patch + 同步 replay.py 的 `_query_pool_snapshot` + analyzer。
+
+### **P2:如果 P0 暴露 polling 是主因,改 polling 实现**
+- 选项 A:把 `/server_info` 改成事件驱动 push(scheduler 在 step 末尾把 stats 写到环形缓冲区,polling 只读不进 scheduler 队列)
+- 选项 B:把 polling 频率从 1Hz 降到 5Hz/10s,在 P1 的拆分数据上验证够用
+- 选项 C:scheduler 端加锁分离,把 stats 读和 admission 决策的临界区拆开
+
+### **P3(条件性,等 P0+P1 数据)**:决定真正的优化方向
+原稿 §7 的 5 条优先级在 `other` 模型纠正后**全部需要重新评估**。等真实拆分数据出来再排。
+
+---
+
+## 8. 局限与 Confounders(已扩充)
+
+1. ⚠️ `held_tokens` 语义在原稿被解读颠倒,引发 `other` 的因果归因错误(已纠正,见 §1.2)。
+2. `other` 字段是计算所得且**未细分**,无法直接归因。需要 P1 instrument 才能区分 radix-cache、running batch、inflight 等。
+3. ⚠️ EXP2+profile 的 415 errors 与 baseline 9 errors **量级差异无法 deconfound**;polling 是 leading suspect 但未证实。**P0 是必经步骤**。
+4. **N=1** 的实验配置:任何 v5+profile vs v5 baseline 的延迟/失败差异都属于 single-run variance 合理范围,**不能作为方向性结论**。
+5. trace 是 single-shot,52 sessions × 4449 reqs 的特定结构可能放大某些路径。
+6. `capacity = 92086` 是 `token_to_kv_pool_allocator.size`,来自 `mem_fraction_static`(未抽具体值),与"H100 80GB 的物理上限"差距是 SGLang 的安全裕量。
+7. ⚠️ §3.1 t=1622 持续高 `other` 30+ tick 的现象 **未与 `last_gen_throughput` 交叉验证**;原稿"running batch + 在途传输"的解释是猜想而非证据。
+8. ⚠️ 18/52 失败 session 的特征(turn_id、input 长度、prefix shape)**未做对比分析**;不能排除某个 session 类型本来就会触发某个固定 bug。
+9. polling 频率 1Hz 错过亚秒级 burst —— `other` 的双峰可能比测到的更剧烈。
+10. critic 指出 `pd-router-d-session-reseed` 在 EXP1 涨(193 vs 152)、EXP2 跌(127 vs 152)的反向移动**未在原稿分析**,这是 admission/路由 决策的清晰信号,应该在 P1 之后回看。
+
+---
+
+## 9. 后续指令(已更新顺序)
+
+1. **P0**: 跑 `scripts/sweep_tp1_v5_baseline_rerun_exp2.sh`,3 次 EXP2 baseline,无 polling。
+2. **P1**: 同时改 SGLang 把 `other` 真正拆开。
+3. 完成 P0+P1 后:
+   - 重跑 EXP2 一次 + 新 instrument(同 polling),拿到 `other` 拆分。
+   - 对比 baseline-rerun 三次的 errors 分布。
+   - 决定是否回退 polling、调 admission、还是攻 specific 18 个 session 的工作负载特征。
+4. 任何 v6 代码改动(优化 admission / eviction / transfer)**必须在 P0+P1 之后**。
+
+---
+
+## 10. 数据产物
+
+```
+outputs/qwen3-30b-tp1-v5-optD-profile/
+├── exp{1,2}_*_metrics.jsonl                # 4449 行 / 实验
+├── exp{1,2}_*_summary.json
+├── exp{1,2}_*_pool_timeseries.jsonl        # 12 MB / 10 MB
+└── kvcache-centric-...20260429T{120847,125911}Z/  # 原始 run dir
+
+outputs/qwen3-30b-tp1-v5-optD/  # baseline 对照(N=1)
+└── exp{1,2}_1p7d_kvc_optD_*
+
+# 待 P0 产生:
+outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/
+└── exp2_2p6d_run{1,2,3}_*
+```
+
+分析脚本:`scripts/analysis/analyze_pool_timeseries.py`(`--json` 拿机器可读输出)。

outputs/qwen3-30b-tp1-v3-kvaware/exp1_1p7d_kvc_kvaware_summary.json

@@ -0,0 +1,88 @@
+{
+  "actual_output_tokens_stats": {
+    "count": 4086.0,
+    "mean": 213.95105237395987,
+    "p50": 83.0,
+    "p90": 562.0,
+    "p99": 1346.0
+  },
+  "cache_hit_request_count": 3929,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 22635.924702180266,
+    "p50": 20010.0,
+    "p90": 48002.0,
+    "p99": 65424.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 363,
+  "execution_modes": {
+    "kvcache-centric": 363,
+    "kvcache-direct-to-d-session": 1716,
+    "pd-router-d-session-reseed": 23,
+    "pd-router-fallback-d-backpressure": 12,
+    "pd-router-fallback-large-append": 5,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 51,
+    "pd-router-fallback-large-append-session-cap": 2148,
+    "pd-router-fallback-no-d-capacity": 7,
+    "pd-router-fallback-session-cap": 32,
+    "pd-router-large-append-reseed": 39,
+    "pd-router-large-append-reseed-after-eviction": 2,
+    "pd-router-turn1-d-backpressure": 1,
+    "pd-router-turn1-no-d-capacity": 3,
+    "pd-router-turn1-seed": 34,
+    "pd-router-turn1-session-cap": 13
+  },
+  "latency_stats_s": {
+    "count": 4086.0,
+    "mean": 4.8753733304192455,
+    "p50": 1.754677688702941,
+    "p90": 12.66968655679375,
+    "p99": 28.717210091650486
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 616,
+    "decode-1": 658,
+    "decode-2": 674,
+    "decode-3": 582,
+    "decode-4": 656,
+    "decode-5": 662,
+    "decode-6": 601
+  },
+  "per_prefill_load": {
+    "prefill-0": 4449
+  },
+  "prefill_request_priorities": {
+    "-100": 98,
+    "100": 2272
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 1716,
+  "total_actual_kv_transfer_blocks": 62123,
+  "total_cached_tokens": 100707229,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4086.0,
+    "mean": 0.005829451223571163,
+    "p50": 0.005684156496173296,
+    "p90": 0.007143743503740225,
+    "p99": 0.008634991403068266
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v3-kvaware/kvcache-centric-kv-aware-worker-admission-20260428T095141Z/sampled-trace.jsonl",
+  "truncated_request_count": 42,
+  "ttft_stats_s": {
+    "count": 4086.0,
+    "mean": 3.5955862397812597,
+    "p50": 0.36274072993546724,
+    "p90": 10.972254231572151,
+    "p99": 27.433656523004174
+  }
+}

outputs/qwen3-30b-tp1-v3-kvaware/exp2_2p6d_kvc_kvaware_summary.json

@@ -0,0 +1,85 @@
+{
+  "actual_output_tokens_stats": {
+    "count": 4440.0,
+    "mean": 225.87972972972972,
+    "p50": 86.0,
+    "p90": 576.0,
+    "p99": 1347.0
+  },
+  "cache_hit_request_count": 4201,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 24345.55787817487,
+    "p50": 21504.0,
+    "p90": 48792.0,
+    "p99": 69120.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 9,
+  "execution_modes": {
+    "kvcache-centric": 9,
+    "kvcache-direct-to-d-session": 1358,
+    "pd-router-d-session-reseed": 12,
+    "pd-router-fallback-d-backpressure": 2,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 52,
+    "pd-router-fallback-large-append-session-cap": 2902,
+    "pd-router-fallback-session-cap": 25,
+    "pd-router-large-append-reseed": 34,
+    "pd-router-large-append-reseed-after-eviction": 4,
+    "pd-router-turn1-d-backpressure": 1,
+    "pd-router-turn1-seed": 30,
+    "pd-router-turn1-session-cap": 20
+  },
+  "latency_stats_s": {
+    "count": 4440.0,
+    "mean": 3.582334662846558,
+    "p50": 1.517257746309042,
+    "p90": 9.225348330102861,
+    "p99": 18.70269925892353
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 710,
+    "decode-1": 630,
+    "decode-2": 763,
+    "decode-3": 737,
+    "decode-4": 879,
+    "decode-5": 730
+  },
+  "per_prefill_load": {
+    "prefill-0": 2225,
+    "prefill-1": 2224
+  },
+  "prefill_request_priorities": {
+    "-100": 80,
+    "100": 3002
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 1358,
+  "total_actual_kv_transfer_blocks": 78979,
+  "total_cached_tokens": 108313387,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4440.0,
+    "mean": 0.005882534704321737,
+    "p50": 0.005807478777200416,
+    "p90": 0.00712956755887717,
+    "p99": 0.008372141476720572
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v3-kvaware/kvcache-centric-kv-aware-worker-admission-20260428T104343Z/sampled-trace.jsonl",
+  "truncated_request_count": 42,
+  "ttft_stats_s": {
+    "count": 4440.0,
+    "mean": 2.2045287611873334,
+    "p50": 0.32809355948120356,
+    "p90": 6.947275545448065,
+    "p99": 16.705802395939827
+  }
+}

outputs/qwen3-30b-tp1-v3-kvaware/sweep_results.txt

@@ -0,0 +1,189 @@
+[2026-04-28 17:51:41] Starting TP1 v3 sweep (KVC with kv-aware policy)
+[2026-04-28 17:51:41] Model: /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+[2026-04-28 17:51:41] Trace: outputs/qwen35-swebench-50sess.jsonl (4449 requests, 52 sessions)
+[2026-04-28 17:51:41] Key change: --policy kv-aware for KVC (was --policy default in v2)
+[2026-04-28 17:51:41] 
+[2026-04-28 17:51:41] === [EXP1] 1P7D KVC kv-aware ===
+[2026-04-28 18:43:43] === exp1_1p7d_kvc_kvaware COMPLETED ===
+[2026-04-28 18:43:43] Summary:
+{
+  "actual_output_tokens_stats": {
+    "count": 4086.0,
+    "mean": 213.95105237395987,
+    "p50": 83.0,
+    "p90": 562.0,
+    "p99": 1346.0
+  },
+  "cache_hit_request_count": 3929,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 22635.924702180266,
+    "p50": 20010.0,
+    "p90": 48002.0,
+    "p99": 65424.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 363,
+  "execution_modes": {
+    "kvcache-centric": 363,
+    "kvcache-direct-to-d-session": 1716,
+    "pd-router-d-session-reseed": 23,
+    "pd-router-fallback-d-backpressure": 12,
+    "pd-router-fallback-large-append": 5,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 51,
+    "pd-router-fallback-large-append-session-cap": 2148,
+    "pd-router-fallback-no-d-capacity": 7,
+    "pd-router-fallback-session-cap": 32,
+    "pd-router-large-append-reseed": 39,
+    "pd-router-large-append-reseed-after-eviction": 2,
+    "pd-router-turn1-d-backpressure": 1,
+    "pd-router-turn1-no-d-capacity": 3,
+    "pd-router-turn1-seed": 34,
+    "pd-router-turn1-session-cap": 13
+  },
+  "latency_stats_s": {
+    "count": 4086.0,
+    "mean": 4.8753733304192455,
+    "p50": 1.754677688702941,
+    "p90": 12.66968655679375,
+    "p99": 28.717210091650486
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 616,
+    "decode-1": 658,
+    "decode-2": 674,
+    "decode-3": 582,
+    "decode-4": 656,
+    "decode-5": 662,
+    "decode-6": 601
+  },
+  "per_prefill_load": {
+    "prefill-0": 4449
+  },
+  "prefill_request_priorities": {
+    "-100": 98,
+    "100": 2272
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 1716,
+  "total_actual_kv_transfer_blocks": 62123,
+  "total_cached_tokens": 100707229,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4086.0,
+    "mean": 0.005829451223571163,
+    "p50": 0.005684156496173296,
+    "p90": 0.007143743503740225,
+    "p99": 0.008634991403068266
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v3-kvaware/kvcache-centric-kv-aware-worker-admission-20260428T095141Z/sampled-trace.jsonl",
+  "truncated_request_count": 42,
+  "ttft_stats_s": {
+    "count": 4086.0,
+    "mean": 3.5955862397812597,
+    "p50": 0.36274072993546724,
+    "p90": 10.972254231572151,
+    "p99": 27.433656523004174
+  }
+}
+[2026-04-28 18:43:43] Saved to outputs/qwen3-30b-tp1-v3-kvaware/exp1_1p7d_kvc_kvaware_summary.json + exp1_1p7d_kvc_kvaware_metrics.jsonl
+[2026-04-28 18:43:43] 
+[2026-04-28 18:43:43] === [EXP2] 2P6D KVC kv-aware ===
+[2026-04-28 19:30:38] === exp2_2p6d_kvc_kvaware COMPLETED ===
+[2026-04-28 19:30:38] Summary:
+{
+  "actual_output_tokens_stats": {
+    "count": 4440.0,
+    "mean": 225.87972972972972,
+    "p50": 86.0,
+    "p90": 576.0,
+    "p99": 1347.0
+  },
+  "cache_hit_request_count": 4201,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 24345.55787817487,
+    "p50": 21504.0,
+    "p90": 48792.0,
+    "p99": 69120.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 9,
+  "execution_modes": {
+    "kvcache-centric": 9,
+    "kvcache-direct-to-d-session": 1358,
+    "pd-router-d-session-reseed": 12,
+    "pd-router-fallback-d-backpressure": 2,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 52,
+    "pd-router-fallback-large-append-session-cap": 2902,
+    "pd-router-fallback-session-cap": 25,
+    "pd-router-large-append-reseed": 34,
+    "pd-router-large-append-reseed-after-eviction": 4,
+    "pd-router-turn1-d-backpressure": 1,
+    "pd-router-turn1-seed": 30,
+    "pd-router-turn1-session-cap": 20
+  },
+  "latency_stats_s": {
+    "count": 4440.0,
+    "mean": 3.582334662846558,
+    "p50": 1.517257746309042,
+    "p90": 9.225348330102861,
+    "p99": 18.70269925892353
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 710,
+    "decode-1": 630,
+    "decode-2": 763,
+    "decode-3": 737,
+    "decode-4": 879,
+    "decode-5": 730
+  },
+  "per_prefill_load": {
+    "prefill-0": 2225,
+    "prefill-1": 2224
+  },
+  "prefill_request_priorities": {
+    "-100": 80,
+    "100": 3002
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 1358,
+  "total_actual_kv_transfer_blocks": 78979,
+  "total_cached_tokens": 108313387,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4440.0,
+    "mean": 0.005882534704321737,
+    "p50": 0.005807478777200416,
+    "p90": 0.00712956755887717,
+    "p99": 0.008372141476720572
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v3-kvaware/kvcache-centric-kv-aware-worker-admission-20260428T104343Z/sampled-trace.jsonl",
+  "truncated_request_count": 42,
+  "ttft_stats_s": {
+    "count": 4440.0,
+    "mean": 2.2045287611873334,
+    "p50": 0.32809355948120356,
+    "p90": 6.947275545448065,
+    "p99": 16.705802395939827
+  }
+}
+[2026-04-28 19:30:38] Saved to outputs/qwen3-30b-tp1-v3-kvaware/exp2_2p6d_kvc_kvaware_summary.json + exp2_2p6d_kvc_kvaware_metrics.jsonl
+[2026-04-28 19:30:38] 
+[2026-04-28 19:30:38] === ALL TP1 V3 SWEEP EXPERIMENTS DONE ===

outputs/qwen3-30b-tp1-v4-cap16/exp1_1p7d_kvc_cap16_summary.json

@@ -0,0 +1,88 @@
+{
+  "actual_output_tokens_stats": {
+    "count": 4014.0,
+    "mean": 215.048081714001,
+    "p50": 83.0,
+    "p90": 570.0,
+    "p99": 1343.0
+  },
+  "cache_hit_request_count": 3865,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 21373.60867610699,
+    "p50": 18429.0,
+    "p90": 45643.0,
+    "p99": 65088.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 435,
+  "execution_modes": {
+    "kvcache-centric": 435,
+    "kvcache-direct-to-d-session": 2180,
+    "pd-router-d-session-reseed": 44,
+    "pd-router-d-session-reseed-after-eviction": 1,
+    "pd-router-fallback-d-backpressure": 36,
+    "pd-router-fallback-large-append": 35,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 52,
+    "pd-router-fallback-large-append-session-cap": 1500,
+    "pd-router-fallback-no-d-capacity": 13,
+    "pd-router-fallback-session-cap": 43,
+    "pd-router-large-append-reseed": 55,
+    "pd-router-large-append-reseed-after-eviction": 3,
+    "pd-router-turn1-d-backpressure": 1,
+    "pd-router-turn1-no-d-capacity": 5,
+    "pd-router-turn1-seed": 46
+  },
+  "latency_stats_s": {
+    "count": 4014.0,
+    "mean": 4.214657033050009,
+    "p50": 1.0827504023909569,
+    "p90": 13.380241627804935,
+    "p99": 24.453291333280504
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 690,
+    "decode-1": 599,
+    "decode-2": 660,
+    "decode-3": 584,
+    "decode-4": 606,
+    "decode-5": 646,
+    "decode-6": 664
+  },
+  "per_prefill_load": {
+    "prefill-0": 4449
+  },
+  "prefill_request_priorities": {
+    "-100": 149,
+    "100": 1685
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 2180,
+  "total_actual_kv_transfer_blocks": 52857,
+  "total_cached_tokens": 95091185,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4014.0,
+    "mean": 0.005804301410418847,
+    "p50": 0.005607025208882987,
+    "p90": 0.007293824862528552,
+    "p99": 0.008864479259402893
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v4-cap16/kvcache-centric-kv-aware-worker-admission-20260428T125022Z/sampled-trace.jsonl",
+  "truncated_request_count": 43,
+  "ttft_stats_s": {
+    "count": 4014.0,
+    "mean": 2.915135478307124,
+    "p50": 0.05643345229327679,
+    "p90": 11.900803190656006,
+    "p99": 22.758968392387033
+  }
+}

outputs/qwen3-30b-tp1-v4-cap16/exp2_2p6d_kvc_cap16_summary.json

@@ -0,0 +1,86 @@
+{
+  "actual_output_tokens_stats": {
+    "count": 4046.0,
+    "mean": 224.65002471576867,
+    "p50": 84.0,
+    "p90": 576.0,
+    "p99": 1349.0
+  },
+  "cache_hit_request_count": 3925,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 22852.7439874129,
+    "p50": 19584.0,
+    "p90": 49009.0,
+    "p99": 67320.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 403,
+  "execution_modes": {
+    "kvcache-centric": 403,
+    "kvcache-direct-to-d-session": 2348,
+    "pd-router-d-session-reseed": 28,
+    "pd-router-fallback-d-backpressure": 7,
+    "pd-router-fallback-large-append": 68,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 45,
+    "pd-router-fallback-large-append-session-cap": 1403,
+    "pd-router-fallback-no-d-capacity": 9,
+    "pd-router-fallback-session-cap": 25,
+    "pd-router-large-append-reseed": 57,
+    "pd-router-large-append-reseed-after-eviction": 6,
+    "pd-router-turn1-no-d-capacity": 1,
+    "pd-router-turn1-seed": 49
+  },
+  "latency_stats_s": {
+    "count": 4046.0,
+    "mean": 2.505981629502371,
+    "p50": 0.8372491216287017,
+    "p90": 6.5139341270551085,
+    "p99": 18.335972285829484
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 767,
+    "decode-1": 680,
+    "decode-2": 906,
+    "decode-3": 818,
+    "decode-4": 800,
+    "decode-5": 478
+  },
+  "per_prefill_load": {
+    "prefill-0": 2225,
+    "prefill-1": 2224
+  },
+  "prefill_request_priorities": {
+    "-100": 140,
+    "100": 1558
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 2348,
+  "total_actual_kv_transfer_blocks": 50727,
+  "total_cached_tokens": 101671858,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4046.0,
+    "mean": 0.005708743129332261,
+    "p50": 0.005565466725497757,
+    "p90": 0.006912594398356141,
+    "p99": 0.008102089307750717
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v4-cap16/kvcache-centric-kv-aware-worker-admission-20260428T134057Z/sampled-trace.jsonl",
+  "truncated_request_count": 36,
+  "ttft_stats_s": {
+    "count": 4046.0,
+    "mean": 1.1653790952959129,
+    "p50": 0.05140436999499798,
+    "p90": 2.6447059931233525,
+    "p99": 15.121314341202378
+  }
+}

outputs/qwen3-30b-tp1-v4-cap16/sweep_results.txt

@@ -0,0 +1,190 @@
+[2026-04-28 20:50:21] Starting TP1 v4 sweep (KVC kv-aware, session soft_cap raised 4->16)
+[2026-04-28 20:50:21] Model: /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+[2026-04-28 20:50:21] Trace: outputs/qwen35-swebench-50sess.jsonl (4449 requests, 52 sessions)
+[2026-04-28 20:50:21] Key change: _decode_session_soft_cap now min(16, ...) instead of min(4, ...)
+[2026-04-28 20:50:21] 
+[2026-04-28 20:50:21] === [EXP1] 1P7D KVC kv-aware cap=16 ===
+[2026-04-28 21:40:57] === exp1_1p7d_kvc_cap16 COMPLETED ===
+[2026-04-28 21:40:57] Summary:
+{
+  "actual_output_tokens_stats": {
+    "count": 4014.0,
+    "mean": 215.048081714001,
+    "p50": 83.0,
+    "p90": 570.0,
+    "p99": 1343.0
+  },
+  "cache_hit_request_count": 3865,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 21373.60867610699,
+    "p50": 18429.0,
+    "p90": 45643.0,
+    "p99": 65088.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 435,
+  "execution_modes": {
+    "kvcache-centric": 435,
+    "kvcache-direct-to-d-session": 2180,
+    "pd-router-d-session-reseed": 44,
+    "pd-router-d-session-reseed-after-eviction": 1,
+    "pd-router-fallback-d-backpressure": 36,
+    "pd-router-fallback-large-append": 35,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 52,
+    "pd-router-fallback-large-append-session-cap": 1500,
+    "pd-router-fallback-no-d-capacity": 13,
+    "pd-router-fallback-session-cap": 43,
+    "pd-router-large-append-reseed": 55,
+    "pd-router-large-append-reseed-after-eviction": 3,
+    "pd-router-turn1-d-backpressure": 1,
+    "pd-router-turn1-no-d-capacity": 5,
+    "pd-router-turn1-seed": 46
+  },
+  "latency_stats_s": {
+    "count": 4014.0,
+    "mean": 4.214657033050009,
+    "p50": 1.0827504023909569,
+    "p90": 13.380241627804935,
+    "p99": 24.453291333280504
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 690,
+    "decode-1": 599,
+    "decode-2": 660,
+    "decode-3": 584,
+    "decode-4": 606,
+    "decode-5": 646,
+    "decode-6": 664
+  },
+  "per_prefill_load": {
+    "prefill-0": 4449
+  },
+  "prefill_request_priorities": {
+    "-100": 149,
+    "100": 1685
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 2180,
+  "total_actual_kv_transfer_blocks": 52857,
+  "total_cached_tokens": 95091185,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4014.0,
+    "mean": 0.005804301410418847,
+    "p50": 0.005607025208882987,
+    "p90": 0.007293824862528552,
+    "p99": 0.008864479259402893
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v4-cap16/kvcache-centric-kv-aware-worker-admission-20260428T125022Z/sampled-trace.jsonl",
+  "truncated_request_count": 43,
+  "ttft_stats_s": {
+    "count": 4014.0,
+    "mean": 2.915135478307124,
+    "p50": 0.05643345229327679,
+    "p90": 11.900803190656006,
+    "p99": 22.758968392387033
+  }
+}
+[2026-04-28 21:40:57] Saved to outputs/qwen3-30b-tp1-v4-cap16/exp1_1p7d_kvc_cap16_summary.json + exp1_1p7d_kvc_cap16_metrics.jsonl
+[2026-04-28 21:40:57] 
+[2026-04-28 21:40:57] === [EXP2] 2P6D KVC kv-aware cap=16 ===
+[2026-04-28 22:27:53] === exp2_2p6d_kvc_cap16 COMPLETED ===
+[2026-04-28 22:27:53] Summary:
+{
+  "actual_output_tokens_stats": {
+    "count": 4046.0,
+    "mean": 224.65002471576867,
+    "p50": 84.0,
+    "p90": 576.0,
+    "p99": 1349.0
+  },
+  "cache_hit_request_count": 3925,
+  "cached_tokens_stats": {
+    "count": 4449.0,
+    "mean": 22852.7439874129,
+    "p50": 19584.0,
+    "p90": 49009.0,
+    "p99": 67320.0
+  },
+  "decode_request_priorities": {},
+  "error_count": 403,
+  "execution_modes": {
+    "kvcache-centric": 403,
+    "kvcache-direct-to-d-session": 2348,
+    "pd-router-d-session-reseed": 28,
+    "pd-router-fallback-d-backpressure": 7,
+    "pd-router-fallback-large-append": 68,
+    "pd-router-fallback-large-append-seed-filter-early-turn": 45,
+    "pd-router-fallback-large-append-session-cap": 1403,
+    "pd-router-fallback-no-d-capacity": 9,
+    "pd-router-fallback-session-cap": 25,
+    "pd-router-large-append-reseed": 57,
+    "pd-router-large-append-reseed-after-eviction": 6,
+    "pd-router-turn1-no-d-capacity": 1,
+    "pd-router-turn1-seed": 49
+  },
+  "latency_stats_s": {
+    "count": 4046.0,
+    "mean": 2.505981629502371,
+    "p50": 0.8372491216287017,
+    "p90": 6.5139341270551085,
+    "p99": 18.335972285829484
+  },
+  "mechanisms": {
+    "kvcache-centric": 4449
+  },
+  "per_decode_load": {
+    "decode-0": 767,
+    "decode-1": 680,
+    "decode-2": 906,
+    "decode-3": 818,
+    "decode-4": 800,
+    "decode-5": 478
+  },
+  "per_prefill_load": {
+    "prefill-0": 2225,
+    "prefill-1": 2224
+  },
+  "prefill_request_priorities": {
+    "-100": 140,
+    "100": 1558
+  },
+  "re_prefill_count": 0,
+  "request_count": 4449,
+  "reuse_expected_count": 4397,
+  "reuse_observed_count": 4397,
+  "router_url": "http://127.0.0.1:8000",
+  "session_reset_count": 0,
+  "session_reused_count": 2348,
+  "total_actual_kv_transfer_blocks": 50727,
+  "total_cached_tokens": 101671858,
+  "total_kv_transfer_blocks": 105235,
+  "tpot_stats_s": {
+    "count": 4046.0,
+    "mean": 0.005708743129332261,
+    "p50": 0.005565466725497757,
+    "p90": 0.006912594398356141,
+    "p99": 0.008102089307750717
+  },
+  "trace_path": "outputs/qwen3-30b-tp1-v4-cap16/kvcache-centric-kv-aware-worker-admission-20260428T134057Z/sampled-trace.jsonl",
+  "truncated_request_count": 36,
+  "ttft_stats_s": {
+    "count": 4046.0,
+    "mean": 1.1653790952959129,
+    "p50": 0.05140436999499798,
+    "p90": 2.6447059931233525,
+    "p99": 15.121314341202378
+  }
+}
+[2026-04-28 22:27:53] Saved to outputs/qwen3-30b-tp1-v4-cap16/exp2_2p6d_kvc_cap16_summary.json + exp2_2p6d_kvc_cap16_metrics.jsonl
+[2026-04-28 22:27:53] 
+[2026-04-28 22:27:53] === ALL TP1 V4 SWEEP EXPERIMENTS DONE ===

pyproject.toml

@@ -0,0 +1,24 @@
+[project]
+name = "agentic-pd-hybrid"
+version = "0.1.0"
+description = "Prototype for session-aware and KV-cache-aware PD routing on SGLang xPyD"
+readme = "README.md"
+requires-python = ">=3.12"
+dependencies = [
+    "httpx>=0.28.1",
+    "mooncake-transfer-engine",
+    "sglang==0.5.10",
+]
+
+[project.scripts]
+agentic-pd-hybrid = "agentic_pd_hybrid.cli:main"
+
+[build-system]
+requires = ["setuptools>=68"]
+build-backend = "setuptools.build_meta"
+
+[tool.setuptools.packages.find]
+where = ["src"]
+
+[tool.uv]
+prerelease = "allow"

scripts/analysis/analyze_backpressure_smoke.py

@@ -0,0 +1,191 @@
+#!/usr/bin/env python3
+"""Analyze backpressure smoke sweep outputs.
+
+For each run dir with a `request-metrics.jsonl` and the new `structural/`
+subdir (admission-events.jsonl, backpressure-events.jsonl,
+session-d-binding.jsonl), report:
+
+- Headline (errors, latency, ttft, direct-to-D rate)
+- Backpressure pause histogram (count, p50/p90 sleep, total pause time per D)
+- Admission probe stats (RPC count, mean RTT, queue_depth distribution,
+  pause_ms distribution)
+- Session pinning (distinct D per session, bimodal direct-to-D rate)
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import statistics
+from collections import Counter, defaultdict
+from pathlib import Path
+
+
+def load_jsonl(path: Path) -> list[dict]:
+    if not path.exists():
+        return []
+    return [json.loads(l) for l in path.open("r", encoding="utf-8") if l.strip()]
+
+
+def summarize_run(run_dir: Path) -> dict:
+    metrics_path = next(run_dir.rglob("request-metrics.jsonl"), None)
+    if metrics_path is None:
+        return {"run_dir": str(run_dir), "error": "no request-metrics.jsonl"}
+
+    summary_path = metrics_path.with_suffix(metrics_path.suffix + ".summary.json")
+    summary = (
+        json.load(summary_path.open()) if summary_path.exists() else {}
+    )
+
+    structural_dir = run_dir / "structural"
+    if not structural_dir.exists():
+        # try metrics dir's parent / structural
+        structural_dir = metrics_path.parent / "structural"
+
+    admission_events = load_jsonl(structural_dir / "admission-events.jsonl")
+    backpressure_events = load_jsonl(structural_dir / "backpressure-events.jsonl")
+    binding_events = load_jsonl(structural_dir / "session-d-binding.jsonl")
+
+    out: dict = {"run_dir": str(run_dir)}
+
+    # Headline metrics from summary.json
+    out["request_count"] = summary.get("request_count")
+    out["error_count"] = summary.get("error_count")
+    out["latency"] = summary.get("latency_stats_s")
+    out["ttft"] = summary.get("ttft_stats_s")
+    out["execution_modes"] = summary.get("execution_modes")
+    out["per_decode_load"] = summary.get("per_decode_load")
+    out["per_prefill_load"] = summary.get("per_prefill_load")
+
+    # Direct-to-D rate from execution_modes
+    em = summary.get("execution_modes", {}) or {}
+    direct = em.get("kvcache-direct-to-d-session", 0)
+    total = sum(em.values()) or 1
+    out["direct_to_d_rate"] = direct / total
+
+    # Session pinning
+    bind_per_session: dict[str, set[int]] = defaultdict(set)
+    for ev in binding_events:
+        bind_per_session[ev["session_id"]].add(ev["decode_worker_index"])
+    if bind_per_session:
+        out["session_count"] = len(bind_per_session)
+        out["avg_distinct_d_per_session"] = (
+            sum(len(v) for v in bind_per_session.values()) / len(bind_per_session)
+        )
+    else:
+        out["session_count"] = 0
+        out["avg_distinct_d_per_session"] = None
+
+    # Direct-to-D rate per session (bimodal check)
+    records = load_jsonl(metrics_path)
+    sess_records: dict[str, list[dict]] = defaultdict(list)
+    for r in records:
+        sess_records[r["session_id"]].append(r)
+    rates = []
+    for sid, turns in sess_records.items():
+        ndir = sum(
+            1 for t in turns if t.get("execution_mode") == "kvcache-direct-to-d-session"
+        )
+        rates.append(ndir / len(turns))
+    if rates:
+        buckets = [0, 0, 0, 0, 0]
+        for r in rates:
+            buckets[min(4, int(r * 5))] += 1
+        out["direct_to_d_rate_buckets"] = {
+            "0-20%": buckets[0],
+            "20-40%": buckets[1],
+            "40-60%": buckets[2],
+            "60-80%": buckets[3],
+            "80-100%": buckets[4],
+        }
+
+    # Backpressure events
+    if backpressure_events:
+        sleeps = [ev["sleep_s"] for ev in backpressure_events]
+        out["backpressure"] = {
+            "event_count": len(backpressure_events),
+            "total_sleep_s": round(sum(sleeps), 2),
+            "sleep_p50_s": round(statistics.median(sleeps), 4),
+            "sleep_p90_s": round(
+                sorted(sleeps)[int(len(sleeps) * 0.9)] if sleeps else 0, 4
+            ),
+            "events_per_d": dict(
+                Counter(ev["server_url"] for ev in backpressure_events).most_common()
+            ),
+        }
+    else:
+        out["backpressure"] = {"event_count": 0, "note": "no backpressure events"}
+
+    # Admission probe stats
+    if admission_events:
+        rtts = [ev["rtt_s"] for ev in admission_events]
+        depths = [ev.get("queue_depth", 0) for ev in admission_events]
+        pauses = [ev.get("recommended_pause_ms", 0) for ev in admission_events]
+        out["admission_probes"] = {
+            "count": len(admission_events),
+            "mean_rtt_s": round(sum(rtts) / len(rtts), 4),
+            "p99_rtt_s": round(sorted(rtts)[int(len(rtts) * 0.99)], 4),
+            "queue_depth_p50": int(statistics.median(depths)),
+            "queue_depth_p90": int(sorted(depths)[int(len(depths) * 0.9)]),
+            "queue_depth_max": max(depths),
+            "pause_ms_p50": int(statistics.median(pauses)),
+            "pause_ms_p90": int(sorted(pauses)[int(len(pauses) * 0.9)]),
+            "pause_ms_max": max(pauses),
+            "nonzero_pause_count": sum(1 for p in pauses if p > 0),
+            "by_reason": dict(
+                Counter(ev.get("reason") or "ok" for ev in admission_events).most_common()
+            ),
+        }
+
+    return out
+
+
+def main() -> None:
+    ap = argparse.ArgumentParser()
+    ap.add_argument("sweep_root", type=Path)
+    ap.add_argument("--json", action="store_true", help="emit JSON only")
+    args = ap.parse_args()
+
+    summaries = []
+    for run_dir in sorted(args.sweep_root.iterdir()):
+        if not run_dir.is_dir():
+            continue
+        summary = summarize_run(run_dir)
+        summaries.append(summary)
+
+    if args.json:
+        print(json.dumps(summaries, indent=2))
+        return
+
+    for s in summaries:
+        print(f"\n{'=' * 70}")
+        print(f"  {s['run_dir']}")
+        print(f"{'=' * 70}")
+        if "error" in s:
+            print(f"  ERROR: {s['error']}")
+            continue
+        print(f"  reqs={s.get('request_count')} errors={s.get('error_count')}")
+        if s.get("latency"):
+            lt = s["latency"]
+            print(
+                f"  latency: mean={lt.get('mean'):.3f} "
+                f"p50={lt.get('p50'):.3f} p90={lt.get('p90'):.3f} p99={lt.get('p99'):.3f}"
+            )
+        if s.get("ttft"):
+            tt = s["ttft"]
+            print(
+                f"  ttft:    mean={tt.get('mean'):.3f} "
+                f"p50={tt.get('p50'):.3f} p90={tt.get('p90'):.3f}"
+            )
+        print(f"  direct_to_d_rate: {s.get('direct_to_d_rate', 0) * 100:.1f}%")
+        print(f"  sessions: {s.get('session_count')} | "
+              f"avg distinct-D-per-session: {s.get('avg_distinct_d_per_session')}")
+        if s.get("direct_to_d_rate_buckets"):
+            print(f"  direct-to-D distribution by session: {s['direct_to_d_rate_buckets']}")
+        if s.get("backpressure"):
+            print(f"  backpressure: {s['backpressure']}")
+        if s.get("admission_probes"):
+            print(f"  admission probes: {s['admission_probes']}")
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/analyze_errors.py

@@ -0,0 +1,83 @@
+#!/usr/bin/env python3
+"""Deep dive into v4 errors: which path, which D, which session, which turn."""
+import json
+import numpy as np
+from pathlib import Path
+from collections import Counter, defaultdict
+
+BASE = Path(__file__).parent
+
+def load_rows(jsonl_path):
+    rows = []
+    with open(jsonl_path) as f:
+        for line in f:
+            rows.append(json.loads(line))
+    return rows
+
+# Compare v3 and v4 errors
+for label, path in [
+    ("v3 1P7D", BASE.parent / "qwen3-30b-tp1-v3-kvaware/exp1_1p7d_kvc_kvaware_metrics.jsonl"),
+    ("v4 1P7D", BASE / "exp1_1p7d_kvc_cap16_metrics.jsonl"),
+    ("v3 2P6D", BASE.parent / "qwen3-30b-tp1-v3-kvaware/exp2_2p6d_kvc_kvaware_metrics.jsonl"),
+    ("v4 2P6D", BASE / "exp2_2p6d_kvc_cap16_metrics.jsonl"),
+]:
+    if not path.exists():
+        print(f"\nSKIP {label}: {path} not found")
+        continue
+    rows = load_rows(path)
+    err = [r for r in rows if r.get("error") is not None]
+    print(f"\n========== {label} ({len(err)} errors / {len(rows)} total = {len(err)/len(rows)*100:.1f}%) ==========")
+
+    # Error finish_reason distribution
+    fr_counter = Counter()
+    for r in err:
+        fr = str(r.get("finish_reason") or r.get("error") or "?")
+        fr_counter[fr[:80]] += 1
+    print(f"finish_reason distribution:")
+    for fr, cnt in fr_counter.most_common():
+        print(f"  {cnt:>4}x  {fr}")
+
+    # Errors by execution mode (these are aborted before mode assignment usually)
+    mode_counter = Counter(r.get("execution_mode", "?") for r in err)
+    print(f"\nerror by execution_mode:")
+    for mode, cnt in mode_counter.most_common():
+        print(f"  {cnt:>4}x  {mode}")
+
+    # Errors per D worker
+    dw_counter = Counter(r.get("assigned_decode_node", "?") for r in err)
+    print(f"\nerror per assigned_decode_node:")
+    for dw, cnt in dw_counter.most_common():
+        print(f"  {cnt:>4}x  {dw}")
+
+    # Errors by turn distribution
+    turn_counter = Counter(r.get("turn_id", -1) for r in err)
+    early = sum(c for t, c in turn_counter.items() if t <= 5)
+    mid = sum(c for t, c in turn_counter.items() if 5 < t <= 30)
+    late = sum(c for t, c in turn_counter.items() if t > 30)
+    print(f"\nerror by turn: early(0-5)={early} mid(6-30)={mid} late(31+)={late}")
+
+    # Per-session error rate
+    per_sess_err = defaultdict(int)
+    per_sess_total = defaultdict(int)
+    for r in rows:
+        per_sess_total[r["session_id"]] += 1
+        if r.get("error") is not None:
+            per_sess_err[r["session_id"]] += 1
+    sess_with_err = [(sid, per_sess_err[sid], per_sess_total[sid]) for sid in per_sess_err]
+    sess_with_err.sort(key=lambda x: -x[1])
+    print(f"\ntop 5 sessions by error count:")
+    for sid, e, t in sess_with_err[:5]:
+        print(f"  session {sid}: {e}/{t} errors ({e/t*100:.0f}%)")
+
+    # Errors timeline: are they bursty?
+    err_ts = sorted([r.get("trace_timestamp_s", 0) for r in err])
+    if err_ts:
+        first_ts = err_ts[0]
+        last_ts = err_ts[-1]
+        all_ts = sorted([r.get("trace_timestamp_s", 0) for r in rows])
+        first_all = all_ts[0]
+        last_all = all_ts[-1]
+        run_duration = last_all - first_all
+        err_first_pct = (err_ts[0] - first_all) / run_duration * 100 if run_duration > 0 else 0
+        err_last_pct = (err_ts[-1] - first_all) / run_duration * 100 if run_duration > 0 else 0
+        print(f"\nerror time range (% of run): {err_first_pct:.1f}% - {err_last_pct:.1f}%")

scripts/analysis/analyze_pool_timeseries.py

@@ -0,0 +1,346 @@
+#!/usr/bin/env python3
+"""Analyze d-pool-timeseries.jsonl produced by --pool-poll-interval-s.
+
+Answers v6's main question: where is D's KV pool actually spent?
+
+For each decode worker, decomposes capacity over the run wall-clock into:
+  - resident_held_active   = held - idle_evictable      (sessions in active use)
+  - resident_held_idle     = idle_evictable             (sessions kept around but evictable)
+  - prefill_backup_or_other = capacity - held - available (everything else: backup blocks,
+                                                          in-flight transfers, fragmentation)
+  - free_available         = available
+
+Also reports session residency churn (how many distinct sessions ever resided per D, and
+how often a session bounced between workers — a strong starvation signal).
+
+Usage:
+  python scripts/analysis/analyze_pool_timeseries.py <run_dir>
+or
+  python scripts/analysis/analyze_pool_timeseries.py <pool_timeseries.jsonl>
+
+Output: human-readable text. Add --json to also print a machine-readable summary.
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import statistics
+from collections import Counter, defaultdict
+from pathlib import Path
+from typing import Any
+
+
+def _load_jsonl(path: Path) -> list[dict[str, Any]]:
+    rows: list[dict[str, Any]] = []
+    with path.open() as fh:
+        for line in fh:
+            line = line.strip()
+            if not line:
+                continue
+            rows.append(json.loads(line))
+    return rows
+
+
+def _resolve_input(path: Path) -> Path:
+    if path.is_file():
+        return path
+    if path.is_dir():
+        candidate = path / "d-pool-timeseries.jsonl"
+        if candidate.is_file():
+            return candidate
+        raise FileNotFoundError(
+            f"{candidate} not found; pass the file directly or a run dir containing it."
+        )
+    raise FileNotFoundError(path)
+
+
+def _percentile(values: list[float], p: float) -> float:
+    if not values:
+        return 0.0
+    s = sorted(values)
+    idx = min(len(s) - 1, max(0, int(round((len(s) - 1) * p))))
+    return s[idx]
+
+
+def _fmt_tokens(n: float) -> str:
+    if n >= 1_000_000:
+        return f"{n / 1_000_000:.2f}M"
+    if n >= 1_000:
+        return f"{n / 1_000:.1f}K"
+    return f"{int(n)}"
+
+
+def _fmt_pct(n: float, total: float) -> str:
+    if total <= 0:
+        return "  -  "
+    return f"{100 * n / total:5.1f}%"
+
+
+def analyze(timeseries_path: Path) -> dict[str, Any]:
+    rows = _load_jsonl(timeseries_path)
+    if not rows:
+        raise ValueError(f"empty timeseries: {timeseries_path}")
+
+    by_worker: dict[str, list[dict[str, Any]]] = defaultdict(list)
+    for row in rows:
+        if row.get("error") and "session_cache_enabled" not in row:
+            # poller failed at this tick — skip
+            continue
+        wid = row.get("worker_id") or "?"
+        by_worker[wid].append(row)
+
+    summary: dict[str, Any] = {
+        "timeseries_path": str(timeseries_path),
+        "total_rows": len(rows),
+        "tick_count": len(by_worker[next(iter(by_worker))]) if by_worker else 0,
+        "wall_s_span": (
+            max(r.get("wall_s", 0.0) for r in rows)
+            - min(r.get("wall_s", 0.0) for r in rows)
+        ),
+        "workers": {},
+    }
+
+    print(f"\n=== Pool timeseries: {timeseries_path}")
+    print(
+        f"  rows={summary['total_rows']}  workers={len(by_worker)}  "
+        f"span={summary['wall_s_span']:.1f}s"
+    )
+
+    # Print per-worker decomposition table
+    header = (
+        f"{'worker':<12} {'role':<8} {'cap':>8} | "
+        f"{'avg_active':>10} {'avg_idle':>10} {'avg_other':>10} {'avg_free':>10} | "
+        f"{'p90_held':>10} {'max_held':>10} {'p90_avail':>10}"
+    )
+    print(header)
+    print("-" * len(header))
+
+    for wid in sorted(by_worker.keys()):
+        ws = by_worker[wid]
+        role = ws[0].get("worker_role", "?")
+        cap_vals = [int(r.get("capacity_tokens") or 0) for r in ws]
+        held_vals = [int(r.get("held_tokens") or 0) for r in ws]
+        avail_vals = [int(r.get("available_tokens") or 0) for r in ws]
+        idle_vals = [int(r.get("idle_evictable_tokens") or 0) for r in ws]
+        # active = held - idle (sessions in active use)
+        active_vals = [max(0, h - i) for h, i in zip(held_vals, idle_vals)]
+        # other = capacity - held - available (prefill backup blocks, in-flight, fragmentation)
+        other_vals = [
+            max(0, c - h - a) for c, h, a in zip(cap_vals, held_vals, avail_vals)
+        ]
+        cap = max(cap_vals) if cap_vals else 0
+
+        avg_active = statistics.fmean(active_vals) if active_vals else 0.0
+        avg_idle = statistics.fmean(idle_vals) if idle_vals else 0.0
+        avg_other = statistics.fmean(other_vals) if other_vals else 0.0
+        avg_avail = statistics.fmean(avail_vals) if avail_vals else 0.0
+
+        p90_held = _percentile([float(v) for v in held_vals], 0.90)
+        max_held = max(held_vals) if held_vals else 0
+        p90_avail = _percentile([float(v) for v in avail_vals], 0.90)
+
+        sess_counts = [int(r.get("session_count") or 0) for r in ws]
+        resident_counts = [int(r.get("resident_session_count") or 0) for r in ws]
+
+        print(
+            f"{wid:<12} {role:<8} {_fmt_tokens(cap):>8} | "
+            f"{_fmt_tokens(avg_active):>4} {_fmt_pct(avg_active, cap):>5} "
+            f"{_fmt_tokens(avg_idle):>4} {_fmt_pct(avg_idle, cap):>5} "
+            f"{_fmt_tokens(avg_other):>4} {_fmt_pct(avg_other, cap):>5} "
+            f"{_fmt_tokens(avg_avail):>4} {_fmt_pct(avg_avail, cap):>5} | "
+            f"{_fmt_tokens(p90_held):>10} {_fmt_tokens(max_held):>10} "
+            f"{_fmt_tokens(p90_avail):>10}"
+        )
+
+        summary["workers"][wid] = {
+            "role": role,
+            "capacity_tokens": cap,
+            "avg_active_held_tokens": avg_active,
+            "avg_idle_evictable_tokens": avg_idle,
+            "avg_other_tokens": avg_other,
+            "avg_available_tokens": avg_avail,
+            "p90_held_tokens": p90_held,
+            "max_held_tokens": max_held,
+            "p90_available_tokens": p90_avail,
+            "max_session_count": max(sess_counts) if sess_counts else 0,
+            "max_resident_session_count": (
+                max(resident_counts) if resident_counts else 0
+            ),
+            "ticks": len(ws),
+        }
+
+    print(
+        "\nLegend: active=held-idle  idle=idle_evictable  "
+        "other=cap-held-avail (radix-protected + running-batch + in-flight + frag)"
+    )
+
+    # P1: decomposition of "other" using pool_breakdown fields (zeros if instrument absent)
+    has_breakdown = any(
+        any(r.get(k) for k in (
+            "radix_evictable_tokens",
+            "radix_protected_tokens",
+            "running_batch_kv_tokens",
+            "transfer_queue_tokens",
+            "prealloc_queue_tokens",
+            "retracted_queue_tokens",
+        ))
+        for r in rows
+    )
+
+    if has_breakdown:
+        print("\n=== P1 'other' decomposition (per worker, mean over run) ===")
+        print(
+            f"{'worker':<12} {'role':<8} | "
+            f"{'r_evictable':>11} {'r_protected':>11} {'slot_private':>12} | "
+            f"{'run_batch':>10} {'transfer':>9} {'prealloc':>9} {'retracted':>10} | "
+            f"{'unaccounted':>11}"
+        )
+        for wid in sorted(by_worker.keys()):
+            ws = by_worker[wid]
+            role = ws[0].get("worker_role", "?")
+            cap = max(int(r.get("capacity_tokens") or 0) for r in ws)
+
+            def m(field: str) -> float:
+                vals = [int(r.get(field) or 0) for r in ws]
+                return statistics.fmean(vals) if vals else 0.0
+
+            r_ev = m("radix_evictable_tokens")
+            r_pr = m("radix_protected_tokens")
+            slot = m("slot_private_held_tokens")
+            rb = m("running_batch_kv_tokens")
+            tq = m("transfer_queue_tokens")
+            pq = m("prealloc_queue_tokens")
+            rq = m("retracted_queue_tokens")
+            avail = m("available_tokens")
+            # `running_batch_kv_tokens` overlaps with radix_protected for tree-tracked
+            # reqs — do NOT subtract it again. Decomposition assumes:
+            # capacity ≈ avail + r_evictable + r_protected + slot_private
+            #           + transfer_queue + prealloc_queue + retracted_queue + unaccounted
+            unacc = max(
+                0,
+                cap - avail - r_ev - r_pr - slot - tq - pq - rq,
+            )
+            print(
+                f"{wid:<12} {role:<8} | "
+                f"{_fmt_tokens(r_ev):>11} {_fmt_tokens(r_pr):>11} {_fmt_tokens(slot):>12} | "
+                f"{_fmt_tokens(rb):>10} {_fmt_tokens(tq):>9} {_fmt_tokens(pq):>9} {_fmt_tokens(rq):>10} | "
+                f"{_fmt_tokens(unacc):>11}"
+            )
+
+            summary["workers"][wid]["pool_breakdown_avg"] = {
+                "radix_evictable": r_ev,
+                "radix_protected": r_pr,
+                "slot_private_held": slot,
+                "running_batch_kv": rb,
+                "transfer_queue": tq,
+                "prealloc_queue": pq,
+                "retracted_queue": rq,
+                "available": avail,
+                "unaccounted": unacc,
+            }
+        print(
+            "\nNote: running_batch_kv_tokens overlaps with radix_protected_tokens "
+            "(tree-tracked decode reqs are also in protected); not summed."
+        )
+    else:
+        print("\n(P1 instrument absent: pool_breakdown fields are all zero)")
+
+    # Session residency churn: how many distinct sessions ever sat on each worker,
+    # and how many sessions hopped across workers (= starvation indicator).
+    print("\n=== Session residency churn ===")
+    sessions_per_worker: dict[str, set[str]] = defaultdict(set)
+    workers_per_session: dict[str, set[str]] = defaultdict(set)
+    resident_ticks_per_session: Counter[str] = Counter()
+    resident_ticks_per_worker: Counter[str] = Counter()
+
+    for row in rows:
+        wid = row.get("worker_id")
+        if wid is None or row.get("worker_role") != "decode":
+            continue
+        sessions = row.get("sessions") or []
+        if not isinstance(sessions, list):
+            continue
+        for entry in sessions:
+            if not isinstance(entry, dict):
+                continue
+            sid = entry.get("session_id")
+            if sid is None:
+                continue
+            if entry.get("resident"):
+                sessions_per_worker[wid].add(sid)
+                workers_per_session[sid].add(wid)
+                resident_ticks_per_session[(wid, sid)] += 1
+                resident_ticks_per_worker[wid] += 1
+
+    # Per-decode worker: distinct session count
+    print(f"  {'worker':<12} {'distinct_sess':>14} {'resident_ticks':>16}")
+    for wid in sorted(sessions_per_worker.keys()):
+        print(
+            f"  {wid:<12} {len(sessions_per_worker[wid]):>14} "
+            f"{resident_ticks_per_worker[wid]:>16}"
+        )
+
+    # Per session: how many workers it hopped across
+    hops = Counter(len(ws) for ws in workers_per_session.values())
+    print(f"\n  Sessions seen on N workers (decode side):")
+    for n, count in sorted(hops.items()):
+        print(f"    on {n} worker(s): {count} sessions")
+
+    starvation = [sid for sid, ws in workers_per_session.items() if len(ws) == 0]
+    multi_hopper = sorted(
+        ((sid, ws) for sid, ws in workers_per_session.items() if len(ws) >= 2),
+        key=lambda x: -len(x[1]),
+    )[:10]
+    if multi_hopper:
+        print(
+            "\n  Top sessions seen resident on multiple workers (potential thrashing):"
+        )
+        for sid, ws in multi_hopper:
+            print(f"    {sid}: {len(ws)} workers ({sorted(ws)})")
+
+    summary["session_residency"] = {
+        "distinct_sessions_per_worker": {
+            wid: len(s) for wid, s in sessions_per_worker.items()
+        },
+        "session_hop_count_distribution": dict(hops),
+        "starvation_session_count": len(starvation),
+    }
+
+    # If a request-metrics file is co-located, also bucket fallback reasons
+    # against contemporaneous pool state (rough — uses tick nearest to median tick).
+    metrics_path = timeseries_path.with_name("request-metrics.jsonl")
+    if metrics_path.exists():
+        print(f"\n=== Request-metrics summary ({metrics_path.name}) ===")
+        mrows = _load_jsonl(metrics_path)
+        modes = Counter(r.get("execution_mode") or "?" for r in mrows)
+        total = sum(modes.values())
+        for mode, count in modes.most_common():
+            print(f"  {count:>6} ({100 * count / total:5.1f}%)  {mode}")
+        summary["execution_modes"] = dict(modes)
+
+    return summary
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument(
+        "path",
+        type=Path,
+        help="Path to d-pool-timeseries.jsonl OR a run dir containing it",
+    )
+    parser.add_argument(
+        "--json",
+        action="store_true",
+        help="Also print a machine-readable JSON summary",
+    )
+    args = parser.parse_args()
+
+    resolved = _resolve_input(args.path)
+    summary = analyze(resolved)
+    if args.json:
+        print("\n=== JSON summary ===")
+        print(json.dumps(summary, indent=2, sort_keys=True, default=str))
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/analyze_ts1_validation.py

@@ -0,0 +1,316 @@
+#!/usr/bin/env python3
+"""TS=1 validation analysis: KVC 1P3D × N=3 + 4DP × 1.
+
+Reads metrics from outputs/qwen3-30b-tp1-ts1-validation/{kvc_1p3d_run{1,2,3},dp4}_metrics.jsonl
+and reports per the structural claims in docs/AGENTIC_FIT_ANALYSIS_ZH.md and TEAM_REPORT.
+
+Sections:
+  1. Headline summary table (errors, latency p50/p90/p99, TTFT p50)
+  2. §1 (session pinning): distinct-D-per-session distribution + direct-to-D bimodal
+  3. §1 (cross-run consistency): sessions consistently starved across all 3 runs + size ratio
+  4. §2 (LRU): KVTransferError counts per D + peak token_usage from worker logs
+  5. §7 (ts=1 vs ts=10): direct-to-D rate, fallback rate, per-D load balance
+  6. KVC vs DP same-scale comparison
+
+Usage: python scripts/analysis/analyze_ts1_validation.py [--root PATH]
+"""
+import argparse
+import json
+import re
+from collections import Counter, defaultdict
+from pathlib import Path
+
+import numpy as np
+
+
+def load_metrics(path):
+    rows = []
+    with open(path) as f:
+        for line in f:
+            line = line.strip()
+            if not line:
+                continue
+            rows.append(json.loads(line))
+    return rows
+
+
+def load_summary(path):
+    with open(path) as f:
+        return json.load(f)
+
+
+def pct(arr, p):
+    if not arr:
+        return float("nan")
+    return float(np.percentile(arr, p))
+
+
+def summarize_run(label, rows, summary):
+    ok = [r for r in rows if r.get("error") is None]
+    err = [r for r in rows if r.get("error") is not None]
+    lats = [r["latency_s"] for r in ok if r.get("latency_s") is not None]
+    ttfts = [r["ttft_s"] for r in ok if r.get("ttft_s") is not None]
+    return {
+        "label": label,
+        "n": len(rows),
+        "ok": len(ok),
+        "err": len(err),
+        "lat_mean": float(np.mean(lats)) if lats else float("nan"),
+        "lat_p50": pct(lats, 50),
+        "lat_p90": pct(lats, 90),
+        "lat_p99": pct(lats, 99),
+        "ttft_mean": float(np.mean(ttfts)) if ttfts else float("nan"),
+        "ttft_p50": pct(ttfts, 50),
+        "summary": summary,
+    }
+
+
+def headline_table(stats):
+    print("\n" + "=" * 110)
+    print("HEADLINE: same trace, same scale, same ts=1")
+    print("=" * 110)
+    cols = ["label", "ok/n", "err", "lat_mean", "lat_p50", "lat_p90", "lat_p99", "ttft_mean", "ttft_p50"]
+    print(f"{cols[0]:<22}{cols[1]:>12}{cols[2]:>6}{cols[3]:>10}{cols[4]:>10}{cols[5]:>10}{cols[6]:>10}{cols[7]:>10}{cols[8]:>10}")
+    for s in stats:
+        ok_n = f"{s['ok']}/{s['n']}"
+        print(f"{s['label']:<22}{ok_n:>12}{s['err']:>6}"
+              f"{s['lat_mean']:>9.3f}s{s['lat_p50']:>9.3f}s{s['lat_p90']:>9.3f}s{s['lat_p99']:>9.3f}s"
+              f"{s['ttft_mean']:>9.3f}s{s['ttft_p50']:>9.3f}s")
+
+
+def session_pinning(rows, label):
+    """§1: distinct D per session — should be ~1.0 if pin behavior persists."""
+    sess_d = defaultdict(set)
+    for r in rows:
+        sid = r.get("session_id")
+        d = r.get("assigned_decode_node") or r.get("decode_node")
+        if sid is not None and d is not None:
+            sess_d[sid].add(d)
+    if not sess_d:
+        return None
+    distinct = [len(s) for s in sess_d.values()]
+    return {
+        "label": label,
+        "n_sessions": len(sess_d),
+        "avg_distinct_D": float(np.mean(distinct)),
+        "max_distinct_D": max(distinct),
+        "sess_d": {sid: sorted(ds) for sid, ds in sess_d.items()},
+    }
+
+
+def direct_to_d_distribution(rows, label):
+    """§1: per-session direct-to-D rate; check for bimodal."""
+    sess_total = Counter()
+    sess_direct = Counter()
+    for r in rows:
+        sid = r.get("session_id")
+        if sid is None:
+            continue
+        sess_total[sid] += 1
+        mode = r.get("execution_mode", "")
+        if mode == "kvcache-direct-to-d-session":
+            sess_direct[sid] += 1
+    rates = []
+    for sid in sess_total:
+        rate = sess_direct[sid] / sess_total[sid]
+        rates.append((sid, rate, sess_total[sid]))
+    bins = [0, 0.2, 0.4, 0.6, 0.8, 1.01]
+    bin_labels = ["0-20%", "20-40%", "40-60%", "60-80%", "80-100%"]
+    counts = [0] * 5
+    for _, r, _ in rates:
+        for i in range(5):
+            if bins[i] <= r < bins[i + 1]:
+                counts[i] += 1
+                break
+    print(f"\n  [{label}] direct-to-D rate distribution (n={len(rates)} sessions):")
+    for lbl, cnt in zip(bin_labels, counts):
+        bar = "█" * cnt
+        print(f"    {lbl:<10}: {cnt:>3}  {bar}")
+    return rates
+
+
+def starved_cross_run(per_run_rates, threshold=0.20):
+    """§1: sessions starved (<threshold direct-to-D) in ALL runs."""
+    if len(per_run_rates) < 2:
+        return None
+    sess_starved = defaultdict(int)
+    sess_lucky = defaultdict(int)
+    for rates in per_run_rates:
+        for sid, rate, _ in rates:
+            if rate < threshold:
+                sess_starved[sid] += 1
+            elif rate > 0.80:
+                sess_lucky[sid] += 1
+    n_runs = len(per_run_rates)
+    consistently_starved = [sid for sid, c in sess_starved.items() if c == n_runs]
+    consistently_lucky = [sid for sid, c in sess_lucky.items() if c == n_runs]
+    return {
+        "n_runs": n_runs,
+        "consistently_starved": consistently_starved,
+        "consistently_lucky": consistently_lucky,
+    }
+
+
+def session_size_comparison(rows, sids_a, sids_b, label_a="A", label_b="B"):
+    """Compare peak input_length of two session groups."""
+    sess_max_input = defaultdict(int)
+    for r in rows:
+        sid = r.get("session_id")
+        ilen = r.get("input_length") or 0
+        if sid is not None and ilen > sess_max_input[sid]:
+            sess_max_input[sid] = ilen
+    a_inputs = [sess_max_input[s] for s in sids_a if s in sess_max_input]
+    b_inputs = [sess_max_input[s] for s in sids_b if s in sess_max_input]
+    if a_inputs and b_inputs:
+        ratio = np.mean(a_inputs) / np.mean(b_inputs)
+        print(f"\n  Cross-run starvation correlates with session size?")
+        print(f"    consistently {label_a} (n={len(a_inputs)}): peak_input mean = {np.mean(a_inputs):.0f}")
+        print(f"    consistently {label_b} (n={len(b_inputs)}): peak_input mean = {np.mean(b_inputs):.0f}")
+        print(f"    {label_a}/{label_b} ratio = {ratio:.2f}x (ts=10 baseline was 1.98x)")
+
+
+def per_d_balance(rows, label):
+    """§7: per-D load balance."""
+    per_d = Counter()
+    for r in rows:
+        d = r.get("assigned_decode_node") or r.get("decode_node")
+        if d:
+            per_d[d] += 1
+    if not per_d:
+        return
+    counts = list(per_d.values())
+    spread = (max(counts) - min(counts)) / max(np.mean(counts), 1)
+    print(f"\n  [{label}] per-D load: {dict(sorted(per_d.items()))}")
+    print(f"    spread (max-min)/mean = {spread*100:.1f}%   "
+          f"(ts=10 KVC 2P6D = ±26%, 8DP CA = ±10%)")
+
+
+def execution_modes_table(rows, label):
+    """Show top execution modes."""
+    ok = [r for r in rows if r.get("error") is None]
+    if not ok:
+        return
+    modes = Counter(r["execution_mode"] for r in ok)
+    print(f"\n  [{label}] execution modes (n_ok={len(ok)}):")
+    for mode, cnt in modes.most_common(8):
+        mode_rows = [r for r in ok if r["execution_mode"] == mode]
+        lats = [r["latency_s"] for r in mode_rows if r.get("latency_s") is not None]
+        ttfts = [r["ttft_s"] for r in mode_rows if r.get("ttft_s") is not None]
+        if lats:
+            print(f"    {mode:<55} {cnt:>5}  ({cnt/len(ok)*100:>4.1f}%)  "
+                  f"lat p50={pct(lats,50):.3f}s p90={pct(lats,90):.3f}s  ttft p50={pct(ttfts,50):.3f}s")
+
+
+def lru_vs_errors(run_dir, label):
+    """§2: trim events vs KVTransferError per worker."""
+    log_dir = run_dir / "logs"
+    if not log_dir.exists():
+        return
+    print(f"\n  [{label}] D-side LRU vs errors (from worker logs):")
+    print(f"    {'worker':<14}{'trim':>8}{'KVTransferError':>20}{'peak_token_usage':>20}")
+    for log_file in sorted(log_dir.glob("decode-*.log")):
+        worker = log_file.stem
+        text = log_file.read_text(errors="ignore")
+        trim_count = len(re.findall(r"Trimmed decode session cache", text))
+        err_count = len(re.findall(r"KVTransferError", text))
+        usages = re.findall(r"token usage: ([\d.]+)", text)
+        peak = max((float(u) for u in usages), default=0.0)
+        print(f"    {worker:<14}{trim_count:>8}{err_count:>20}{peak:>20.3f}")
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--root", default="outputs/qwen3-30b-tp1-ts1-validation",
+                        help="Sweep output root")
+    args = parser.parse_args()
+
+    root = Path(args.root)
+    if not root.is_absolute():
+        root = Path("/mnt/kzlin/workflow/pd-hybrid/agentic-pd-hybrid") / root
+
+    # Load all available runs
+    stats = []
+    rows_by_run = {}
+    for label in ("kvc_1p3d_run1", "kvc_1p3d_run2", "kvc_1p3d_run3", "dp4"):
+        m = root / f"{label}_metrics.jsonl"
+        s = root / f"{label}_summary.json"
+        if not m.exists() or not s.exists():
+            print(f"  [{label}] not yet available ({m.name})")
+            continue
+        rows = load_metrics(m)
+        summary = load_summary(s)
+        rows_by_run[label] = rows
+        stats.append(summarize_run(label, rows, summary))
+
+    if not stats:
+        print("No runs available yet.")
+        return
+
+    # 1. Headline table
+    headline_table(stats)
+
+    # 2. §1 session pinning per KVC run + per-D balance + execution modes
+    print("\n" + "=" * 110)
+    print("§1 / §7: SESSION PINNING + LOAD BALANCE")
+    print("=" * 110)
+    per_run_rates = []
+    for label, rows in rows_by_run.items():
+        if not label.startswith("kvc_"):
+            continue
+        pin = session_pinning(rows, label)
+        if pin:
+            print(f"\n  [{label}] sessions={pin['n_sessions']}  "
+                  f"avg_distinct_D={pin['avg_distinct_D']:.2f}  "
+                  f"max_distinct_D={pin['max_distinct_D']}  "
+                  f"(ts=10 baseline avg=1.00 → 100% pin)")
+        rates = direct_to_d_distribution(rows, label)
+        per_run_rates.append(rates)
+        per_d_balance(rows, label)
+        execution_modes_table(rows, label)
+
+    # 3. §1 cross-run starvation
+    if len(per_run_rates) >= 2:
+        print("\n" + "=" * 110)
+        print(f"§1 CROSS-RUN STARVATION (across {len(per_run_rates)} KVC runs)")
+        print("=" * 110)
+        cross = starved_cross_run(per_run_rates)
+        if cross:
+            n_starved = len(cross["consistently_starved"])
+            n_lucky = len(cross["consistently_lucky"])
+            print(f"\n  Sessions starved (<20% direct-to-D) in all {cross['n_runs']} runs: {n_starved}")
+            print(f"  Sessions lucky (>80% direct-to-D) in all {cross['n_runs']} runs: {n_lucky}")
+            print(f"  (ts=10 baseline: 13/52 starved, 14/52 lucky — extreme bimodal)")
+            # session size comparison from run 1
+            if "kvc_1p3d_run1" in rows_by_run and n_starved and n_lucky:
+                session_size_comparison(rows_by_run["kvc_1p3d_run1"],
+                                        cross["consistently_starved"],
+                                        cross["consistently_lucky"],
+                                        "starved", "lucky")
+
+    # 4. §2 D-side LRU vs errors from raw logs
+    print("\n" + "=" * 110)
+    print("§2: D-SIDE LRU TRIM vs KVTransferError (from worker logs)")
+    print("=" * 110)
+    for label in rows_by_run:
+        if not label.startswith("kvc_"):
+            continue
+        # find the matching raw run dir
+        run_dirs = sorted(root.glob("kvcache-centric-*/"))
+        if not run_dirs:
+            continue
+        # naive: index matches run order; could be wrong if dirs got reordered
+        idx = int(label.split("run")[-1]) - 1
+        if idx < len(run_dirs):
+            lru_vs_errors(run_dirs[idx], label)
+
+    # 5. DP-only inspection
+    if "dp4" in rows_by_run:
+        print("\n" + "=" * 110)
+        print("4DP CA SANITY")
+        print("=" * 110)
+        per_d_balance(rows_by_run["dp4"], "dp4")
+        execution_modes_table(rows_by_run["dp4"], "dp4")
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/analyze_v3.py

@@ -0,0 +1,89 @@
+#!/usr/bin/env python3
+"""Analyze v3 (kv-aware) results — find why fallback-large-append-session-cap dominates."""
+import json
+import numpy as np
+from pathlib import Path
+from collections import Counter, defaultdict
+
+BASE = Path(__file__).parent
+
+def load_rows(jsonl_path):
+    rows = []
+    with open(jsonl_path) as f:
+        for line in f:
+            rows.append(json.loads(line))
+    return rows
+
+exp1 = load_rows(BASE / "exp1_1p7d_kvc_kvaware_metrics.jsonl")
+exp2 = load_rows(BASE / "exp2_2p6d_kvc_kvaware_metrics.jsonl")
+
+for name, rows in [("Exp1 1P7D", exp1), ("Exp2 2P6D", exp2)]:
+    print(f"\n========== {name} ==========")
+    ok = [r for r in rows if r.get("error") is None]
+
+    # Execution mode breakdown by latency
+    modes = Counter(r["execution_mode"] for r in ok)
+    print(f"\nExecution modes (n={len(ok)}):")
+    for mode, count in modes.most_common():
+        mode_rows = [r for r in ok if r["execution_mode"] == mode]
+        lats = [r["latency_s"] for r in mode_rows]
+        ttfts = [r["ttft_s"] for r in mode_rows]
+        print(f"  {mode}: n={count} ({count/len(ok)*100:.1f}%) "
+              f"lat P50={np.percentile(lats,50):.3f}s P90={np.percentile(lats,90):.3f}s | "
+              f"ttft P50={np.percentile(ttfts,50):.3f}s P90={np.percentile(ttfts,90):.3f}s")
+
+    # Per-D session distribution
+    per_d_sessions = defaultdict(set)
+    for r in ok:
+        d = r.get("assigned_decode_node", "?")
+        per_d_sessions[d].add(r["session_id"])
+    print(f"\nSessions per D worker:")
+    for d in sorted(per_d_sessions.keys()):
+        print(f"  {d}: {len(per_d_sessions[d])} unique sessions")
+
+    # session-cap fallback analysis
+    sc_rows = [r for r in ok if r["execution_mode"] == "pd-router-fallback-large-append-session-cap"]
+    if sc_rows:
+        print(f"\nSession-cap fallback details (n={len(sc_rows)}):")
+        # Which sessions hit this most?
+        sc_per_sess = Counter(r["session_id"] for r in sc_rows)
+        print(f"  Sessions hitting session-cap (top 5):")
+        for sid, cnt in sc_per_sess.most_common(5):
+            print(f"    session {sid}: {cnt} times")
+        # Per-D distribution
+        sc_per_d = Counter(r.get("assigned_decode_node", "?") for r in sc_rows)
+        print(f"  Per-D distribution: {dict(sc_per_d.most_common())}")
+        # Input length distribution
+        inp = [r.get("input_length", 0) for r in sc_rows]
+        print(f"  Input length: P50={np.percentile(inp,50):.0f} P90={np.percentile(inp,90):.0f}")
+        # Turn distribution
+        turns = Counter(r.get("turn_id", -1) for r in sc_rows)
+        print(f"  Turn distribution (top 5): {dict(turns.most_common(5))}")
+
+    # Direct-to-D analysis (ideal path)
+    dd_rows = [r for r in ok if r["execution_mode"] == "kvcache-direct-to-d-session"]
+    if dd_rows:
+        lats = [r["latency_s"] for r in dd_rows]
+        ttfts = [r["ttft_s"] for r in dd_rows]
+        kv_blocks = [r.get("actual_kv_transfer_blocks", 0) for r in dd_rows]
+        cached = [r.get("cached_tokens", 0) for r in dd_rows]
+        print(f"\nDirect-to-D details (n={len(dd_rows)}):")
+        print(f"  lat P50={np.percentile(lats,50):.3f}s P90={np.percentile(lats,90):.3f}s P99={np.percentile(lats,99):.3f}s")
+        print(f"  ttft P50={np.percentile(ttfts,50):.3f}s P90={np.percentile(ttfts,90):.3f}s")
+        print(f"  KV transfer: P50={np.percentile(kv_blocks,50):.0f} (should be 0 — no P involved)")
+        print(f"  cached_tokens P50={np.percentile(cached,50):.0f}")
+
+    # Sessions: how many turns each, how many used direct-to-d
+    print(f"\nPer-session direct-to-D rate (top 10 by total turns):")
+    per_sess = defaultdict(list)
+    for r in ok:
+        per_sess[r["session_id"]].append(r)
+    sess_stats = []
+    for sid, sreqs in per_sess.items():
+        total = len(sreqs)
+        dd = sum(1 for r in sreqs if r["execution_mode"] == "kvcache-direct-to-d-session")
+        sc = sum(1 for r in sreqs if "session-cap" in r["execution_mode"])
+        sess_stats.append((sid, total, dd, sc))
+    sess_stats.sort(key=lambda x: -x[1])
+    for sid, total, dd, sc in sess_stats[:10]:
+        print(f"  session {sid}: {total} turns, {dd} direct-to-D ({dd/total*100:.0f}%), {sc} session-cap fallback ({sc/total*100:.0f}%)")

scripts/analysis/analyze_v4.py

@@ -0,0 +1,52 @@
+#!/usr/bin/env python3
+"""V4 results analysis: errors, execution modes, latency by mode."""
+import json
+import numpy as np
+from pathlib import Path
+from collections import Counter
+
+BASE = Path(__file__).parent
+
+def load_rows(jsonl_path):
+    rows = []
+    with open(jsonl_path) as f:
+        for line in f:
+            rows.append(json.loads(line))
+    return rows
+
+for name, path in [
+    ("Exp1 1P7D cap=16", BASE / "exp1_1p7d_kvc_cap16_metrics.jsonl"),
+    ("Exp2 2P6D cap=16", BASE / "exp2_2p6d_kvc_cap16_metrics.jsonl"),
+]:
+    rows = load_rows(path)
+    print(f"\n========== {name} ==========")
+    ok = [r for r in rows if r.get("error") is None]
+    err = [r for r in rows if r.get("error") is not None]
+    print(f"Total: {len(rows)}, OK: {len(ok)}, Errors: {len(err)}")
+
+    # Errors finish_reason
+    if err:
+        finish_reasons = Counter()
+        for r in err:
+            fr = str(r.get("finish_reason") or r.get("error") or "?")
+            # Truncate long messages
+            short = fr[:120]
+            finish_reasons[short] += 1
+        print(f"\nError finish_reasons (top 5):")
+        for fr, cnt in finish_reasons.most_common(5):
+            print(f"  {cnt}x: {fr}")
+
+    # Execution mode latency breakdown
+    modes = Counter(r["execution_mode"] for r in ok)
+    print(f"\nTop execution modes by latency:")
+    print(f"{'mode':<55}{'n':<8}{'%':<8}{'P50 lat':<10}{'P90 lat':<10}{'TTFT P50':<10}")
+    for mode, count in modes.most_common(8):
+        mode_rows = [r for r in ok if r["execution_mode"] == mode]
+        lats = [r["latency_s"] for r in mode_rows]
+        ttfts = [r["ttft_s"] for r in mode_rows]
+        print(f"  {mode:<53}{count:<8}{count/len(ok)*100:>5.1f}%  {np.percentile(lats,50):>7.3f}s  {np.percentile(lats,90):>7.3f}s  {np.percentile(ttfts,50):>7.3f}s")
+
+    # Per-D load
+    per_d = Counter(r.get("assigned_decode_node", "?") for r in ok)
+    print(f"\nPer-D load: max/min ratio = {max(per_d.values())/max(min(per_d.values()),1):.2f}x")
+    print(f"  {dict(per_d.most_common())}")

scripts/analysis/compare_no_error.py

@@ -0,0 +1,136 @@
+#!/usr/bin/env python3
+"""Compare KVC variants vs baseline, EXCLUDING errors and truncated requests."""
+import json
+import numpy as np
+from pathlib import Path
+
+OUT = Path("/mnt/kzlin/workflow/pd-hybrid/agentic-pd-hybrid/outputs")
+
+DATASETS = [
+    ("baseline 8DP", OUT / "qwen3-30b-tp1-v2-fixed/exp1_8way_dp_cache_aware_metrics.jsonl"),
+    ("v3 1P7D",     OUT / "qwen3-30b-tp1-v3-kvaware/exp1_1p7d_kvc_kvaware_metrics.jsonl"),
+    ("v3 2P6D",     OUT / "qwen3-30b-tp1-v3-kvaware/exp2_2p6d_kvc_kvaware_metrics.jsonl"),
+    ("v4 1P7D",     OUT / "qwen3-30b-tp1-v4-cap16/exp1_1p7d_kvc_cap16_metrics.jsonl"),
+    ("v4 2P6D",     OUT / "qwen3-30b-tp1-v4-cap16/exp2_2p6d_kvc_cap16_metrics.jsonl"),
+]
+
+def load_rows(path):
+    rows = []
+    with open(path) as f:
+        for line in f:
+            rows.append(json.loads(line))
+    return rows
+
+def is_truncated(row):
+    a = row.get("actual_output_tokens")
+    r = row.get("requested_output_tokens")
+    if a is not None and r is not None and r > 1:
+        return a < r * 0.5
+    return False
+
+def stats(values):
+    if not values:
+        return {"n": 0}
+    a = np.array(values)
+    return {
+        "n":    len(a),
+        "mean": float(np.mean(a)),
+        "p50":  float(np.percentile(a, 50)),
+        "p90":  float(np.percentile(a, 90)),
+        "p99":  float(np.percentile(a, 99)),
+    }
+
+def fmt(s, key):
+    if s["n"] == 0:
+        return "N/A"
+    v = s[key]
+    return f"{v:.3f}s" if v < 100 else f"{v:.1f}s"
+
+results = []
+for label, path in DATASETS:
+    if not path.exists():
+        print(f"SKIP {label}")
+        continue
+    rows = load_rows(path)
+    total = len(rows)
+    err_n = sum(1 for r in rows if r.get("error") is not None)
+    trunc_n = sum(1 for r in rows if r.get("error") is None and is_truncated(r))
+
+    # Filter: error=None AND not truncated AND latency present
+    clean = [r for r in rows
+             if r.get("error") is None
+             and not is_truncated(r)
+             and r.get("latency_s") is not None]
+
+    lats = [r["latency_s"] for r in clean]
+    ttfts = [r["ttft_s"] for r in clean if r.get("ttft_s") is not None]
+
+    results.append({
+        "label": label,
+        "total": total,
+        "err": err_n,
+        "trunc": trunc_n,
+        "clean_n": len(clean),
+        "lat": stats(lats),
+        "ttft": stats(ttfts),
+    })
+
+# Print comparison table
+print(f"\n{'='*100}")
+print("LATENCY (excluding errors AND truncated)")
+print(f"{'='*100}")
+print(f"{'config':<16}{'total':>7}{'err':>6}{'trunc':>7}{'clean':>7}  {'mean':>9}{'P50':>9}{'P90':>9}{'P99':>9}")
+for r in results:
+    print(f"{r['label']:<16}{r['total']:>7}{r['err']:>6}{r['trunc']:>7}{r['clean_n']:>7}  "
+          f"{fmt(r['lat'],'mean'):>9}{fmt(r['lat'],'p50'):>9}{fmt(r['lat'],'p90'):>9}{fmt(r['lat'],'p99'):>9}")
+
+print(f"\n{'='*100}")
+print("TTFT (excluding errors AND truncated)")
+print(f"{'='*100}")
+print(f"{'config':<16}{'clean':>7}  {'mean':>9}{'P50':>9}{'P90':>9}{'P99':>9}")
+for r in results:
+    print(f"{r['label']:<16}{r['clean_n']:>7}  "
+          f"{fmt(r['ttft'],'mean'):>9}{fmt(r['ttft'],'p50'):>9}{fmt(r['ttft'],'p90'):>9}{fmt(r['ttft'],'p99'):>9}")
+
+# Also: per-execution-mode breakdown for v4 only (the most interesting)
+print(f"\n{'='*100}")
+print("V4 2P6D: per-execution-mode (excluding errors and truncated)")
+print(f"{'='*100}")
+v4_2p6d = next((p for l, p in DATASETS if l == "v4 2P6D"), None)
+if v4_2p6d:
+    rows = load_rows(v4_2p6d)
+    clean = [r for r in rows if r.get("error") is None and not is_truncated(r)]
+    from collections import Counter
+    modes = Counter(r["execution_mode"] for r in clean)
+    print(f"{'mode':<55}{'n':>7}{'%':>7}  {'mean':>9}{'P50':>9}{'P90':>9}{'P99':>9}")
+    for mode, count in modes.most_common(10):
+        m_rows = [r for r in clean if r["execution_mode"] == mode]
+        s = stats([r["latency_s"] for r in m_rows])
+        pct = count/len(clean)*100
+        print(f"  {mode:<53}{count:>7}{pct:>6.1f}%  {fmt(s,'mean'):>9}{fmt(s,'p50'):>9}{fmt(s,'p90'):>9}{fmt(s,'p99'):>9}")
+
+# Also: WHAT IF we only count direct-to-D? (Pure KVC performance)
+print(f"\n{'='*100}")
+print("Pure KVC (kvcache-direct-to-d-session ONLY) vs Baseline")
+print(f"{'='*100}")
+for label, path in DATASETS:
+    if not path.exists() or "1P7D" not in label and "2P6D" not in label:
+        continue
+    rows = load_rows(path)
+    direct = [r for r in rows
+              if r.get("error") is None and not is_truncated(r)
+              and r.get("execution_mode") == "kvcache-direct-to-d-session"]
+    if not direct:
+        continue
+    s_lat = stats([r["latency_s"] for r in direct])
+    s_ttft = stats([r["ttft_s"] for r in direct if r.get("ttft_s") is not None])
+    print(f"{label:<16}n={s_lat['n']:>5}  lat: P50={fmt(s_lat,'p50')} P90={fmt(s_lat,'p90')}  ttft: P50={fmt(s_ttft,'p50')} P90={fmt(s_ttft,'p90')}")
+
+# Baseline for reference (already non-fallback by definition)
+print()
+baseline_path = OUT / "qwen3-30b-tp1-v2-fixed/exp1_8way_dp_cache_aware_metrics.jsonl"
+baseline_rows = load_rows(baseline_path)
+clean = [r for r in baseline_rows if r.get("error") is None and not is_truncated(r)]
+s_lat = stats([r["latency_s"] for r in clean])
+s_ttft = stats([r["ttft_s"] for r in clean if r.get("ttft_s") is not None])
+print(f"{'baseline 8DP':<16}n={s_lat['n']:>5}  lat: P50={fmt(s_lat,'p50')} P90={fmt(s_lat,'p90')}  ttft: P50={fmt(s_ttft,'p50')} P90={fmt(s_ttft,'p90')}")

scripts/analysis/plot_cache_efficiency.py

@@ -0,0 +1,209 @@
+#!/usr/bin/env python3
+"""Cache efficiency comparison: KVC 1P3D v2 vs 4-way DP CA.
+
+Generates docs/figures/cache_efficiency.png — two-panel:
+  left:  cache hit rate vs turn number   (mechanism: affinity vs LRU)
+  right: ECDF of per-request uncached tokens  (per-request impact)
+
+Resolves the apparent paradox: KVC has 27% less total KV pool capacity
+(3 × 92K = 276K  vs  DP 4 × 87K = 351K) yet achieves higher cache hit rate
+(98.1% vs 96.8%) and lower mean uncached tokens per request (560 vs 952).
+
+The left panel shows the mechanism: KVC's session affinity makes cache hit
+rate grow with turn count (more cache accumulates on the pinned D), while
+DP's hash + radix-LRU causes cache hit rate to decay through the middle
+turns (other sessions' KV competes via LRU eviction).
+
+The right panel quantifies the impact: KVC's uncached tokens are
+concentrated near 0 (mean 560), DP's are spread (mean 952).
+
+Aborted / errored requests are excluded.
+"""
+
+from __future__ import annotations
+
+import json
+from collections import defaultdict
+from pathlib import Path
+
+import matplotlib.pyplot as plt
+import numpy as np
+
+ROOT = Path(__file__).resolve().parents[2]
+KVC = ROOT / "outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl"
+DP = ROOT / "outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl"
+OUT = ROOT / "docs/figures/cache_efficiency.png"
+
+
+def load(p: Path) -> list[dict]:
+    return [json.loads(line) for line in p.open()]
+
+
+def is_failed(r: dict) -> bool:
+    if r.get("error"):
+        return True
+    fr = r.get("finish_reason")
+    if fr and ("abort" in str(fr).lower() or "badrequest" in str(fr).lower()):
+        return True
+    return False
+
+
+def main() -> None:
+    kvc = [r for r in load(KVC) if not is_failed(r)]
+    dp = [r for r in load(DP) if not is_failed(r)]
+
+    KVC_COLOR = "#1F77B4"
+    DP_COLOR = "#D62728"
+
+    fig, axes = plt.subplots(1, 2, figsize=(15, 6.5))
+
+    # ------------------------------------------------------------------
+    # Left panel: cache hit rate per turn
+    # Bin requests by turn_id, plot mean hit rate per bin with shaded band
+    # ------------------------------------------------------------------
+    def bin_by_turn(rows: list[dict]) -> tuple[list[int], list[float], list[float], list[float]]:
+        per_turn: defaultdict[int, list[float]] = defaultdict(list)
+        for r in rows:
+            if r["input_length"] == 0:
+                continue
+            hit = r.get("cached_tokens", 0) / r["input_length"]
+            per_turn[r["turn_id"]].append(hit)
+        turns = sorted(per_turn.keys())
+        means, p25s, p75s = [], [], []
+        for t in turns:
+            arr = np.array(per_turn[t])
+            means.append(float(np.mean(arr)))
+            p25s.append(float(np.quantile(arr, 0.25)))
+            p75s.append(float(np.quantile(arr, 0.75)))
+        return turns, means, p25s, p75s
+
+    kvc_t, kvc_m, kvc_lo, kvc_hi = bin_by_turn(kvc)
+    dp_t, dp_m, dp_lo, dp_hi = bin_by_turn(dp)
+
+    # Cap x-axis: tails get noisy below ~5 samples per bin
+    max_turn = 100
+
+    ax = axes[0]
+    ax.plot(kvc_t, kvc_m, color=KVC_COLOR, lw=2.5,
+            label=f"KVC 1P3D v2  (overall hit 98.1%)")
+    ax.fill_between(kvc_t, kvc_lo, kvc_hi, color=KVC_COLOR, alpha=0.18,
+                    label="KVC IQR (p25-p75)")
+    ax.plot(dp_t, dp_m, color=DP_COLOR, lw=2.5,
+            label=f"4-way DP CA  (overall hit 96.8%)")
+    ax.fill_between(dp_t, dp_lo, dp_hi, color=DP_COLOR, alpha=0.18,
+                    label="DP IQR (p25-p75)")
+
+    # Annotate the mid-turn drift gap
+    drift_turns = list(range(8, 25))
+    drift_kvc = np.mean([m for t, m in zip(kvc_t, kvc_m) if t in drift_turns])
+    drift_dp = np.mean([m for t, m in zip(dp_t, dp_m) if t in drift_turns])
+    ax.axvspan(8, 25, color="#999", alpha=0.08, label="_nolegend_")
+    ax.text(16, 0.65,
+            f"Mid-turn region\n(turns 8-25):\nKVC {drift_kvc*100:.1f}%  |  DP {drift_dp*100:.1f}%\nGap {(drift_kvc-drift_dp)*100:+.1f} pp",
+            ha="center", va="center", fontsize=9.5,
+            bbox=dict(facecolor="white", edgecolor="gray", alpha=0.92, pad=4))
+
+    ax.set_xlim(1, max_turn)
+    ax.set_ylim(0.4, 1.02)
+    ax.set_xlabel("Turn number within session", fontsize=11)
+    ax.set_ylabel("Per-request cache hit rate (cached / input_length)", fontsize=11)
+    ax.set_title("Cache hit rate vs turn number\n(mechanism: session affinity vs hash-LRU)",
+                 fontsize=12, pad=10)
+    ax.legend(loc="lower right", fontsize=9.5, framealpha=0.95)
+    ax.grid(True, linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+
+    # ------------------------------------------------------------------
+    # Right panel: ECDF of per-request uncached tokens (log x)
+    # ------------------------------------------------------------------
+    def ecdf(rows: list[dict]) -> tuple[np.ndarray, np.ndarray]:
+        vals = np.array([
+            max(1, r["input_length"] - r.get("cached_tokens", 0))
+            for r in rows
+        ])
+        vals = np.sort(vals)
+        return vals, np.arange(1, len(vals) + 1) / len(vals)
+
+    kvc_x, kvc_y = ecdf(kvc)
+    dp_x, dp_y = ecdf(dp)
+
+    ax = axes[1]
+    ax.plot(kvc_x, kvc_y, color=KVC_COLOR, lw=2.5,
+            label=f"KVC 1P3D v2  (mean {int(np.mean(kvc_x))} tokens)")
+    ax.plot(dp_x, dp_y, color=DP_COLOR, lw=2.5,
+            label=f"4-way DP CA  (mean {int(np.mean(dp_x))} tokens)")
+
+    # Median markers
+    kvc_p50 = np.quantile(kvc_x, 0.50)
+    dp_p50 = np.quantile(dp_x, 0.50)
+    ax.axhline(0.5, color="gray", linestyle=":", alpha=0.5)
+    ax.text(1.2, 0.52, "median (50% of requests below this)",
+            fontsize=8.5, color="gray", style="italic")
+    ax.axvline(kvc_p50, color=KVC_COLOR, ls="--", alpha=0.5, lw=1.0)
+    ax.axvline(dp_p50, color=DP_COLOR, ls="--", alpha=0.5, lw=1.0)
+    ax.text(kvc_p50, 0.06, f"KVC\nmedian\n{int(kvc_p50)}",
+            color=KVC_COLOR, fontsize=9, ha="center", va="bottom",
+            bbox=dict(facecolor="white", edgecolor="none", alpha=0.75, pad=1))
+    ax.text(dp_p50, 0.06, f"DP\nmedian\n{int(dp_p50)}",
+            color=DP_COLOR, fontsize=9, ha="center", va="bottom",
+            bbox=dict(facecolor="white", edgecolor="none", alpha=0.75, pad=1))
+
+    # Annotate the separation: at uncached = 500 tokens, what fraction below?
+    sep_x = 500
+    kvc_at_sep = (kvc_x <= sep_x).mean()
+    dp_at_sep = (dp_x <= sep_x).mean()
+    ax.axvline(sep_x, color="#666", linestyle=":", alpha=0.6, lw=1.0)
+    ax.annotate(
+        f"At uncached = {sep_x} tokens:\n"
+        f"KVC {kvc_at_sep*100:.0f}% of requests below\n"
+        f"DP  {dp_at_sep*100:.0f}% of requests below",
+        xy=(sep_x, dp_at_sep),
+        xytext=(2500, 0.35),
+        fontsize=9.5,
+        bbox=dict(facecolor="white", edgecolor="gray", alpha=0.92, pad=4),
+        arrowprops=dict(arrowstyle="->", color="#666", lw=0.8),
+    )
+
+    ax.set_xscale("log")
+    ax.set_xlim(1, 1e5)
+    ax.set_xticks([1, 10, 100, 1000, 10000, 100000])
+    ax.set_xticklabels(["1", "10", "100", "1K", "10K", "100K"])
+    ax.set_ylim(0, 1.02)
+    ax.set_xlabel("Uncached tokens per request  (log scale)", fontsize=11)
+    ax.set_ylabel("Cumulative fraction of requests", fontsize=11)
+    ax.set_title("ECDF of uncached tokens per request\n(impact: KVC concentrates near zero)",
+                 fontsize=12, pad=10)
+    ax.legend(loc="lower right", fontsize=10, framealpha=0.95)
+    ax.grid(True, which="both", linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+
+    fig.suptitle(
+        "Cache efficiency paradox:  KVC has 27% LESS total KV pool (276K vs 351K tokens) yet caches MORE per request.\n"
+        "Left: session-affinity lets KVC's cache accumulate with turns; DP's hash-LRU loses cache to cross-session competition.\n"
+        "Right: net effect — KVC's uncached compute is concentrated near zero, DP's is spread over 100-10K tokens.",
+        fontsize=11.5, y=1.05,
+    )
+    plt.tight_layout()
+    plt.savefig(OUT, dpi=150, bbox_inches="tight")
+    print(f"wrote {OUT}")
+    plt.close(fig)
+
+    # ------------------------------------------------------------------
+    # Print summary for doc reference
+    # ------------------------------------------------------------------
+    print("\n=== Cache efficiency stats ===")
+    print(f"KVC v2:  total_input={sum(r['input_length'] for r in kvc)/1e6:.1f}M tokens")
+    print(f"         total_cached={sum(r.get('cached_tokens',0) for r in kvc)/1e6:.1f}M tokens")
+    print(f"         hit rate {sum(r.get('cached_tokens',0) for r in kvc)/sum(r['input_length'] for r in kvc)*100:.2f}%")
+    print(f"         mean uncached {np.mean(kvc_x):.0f}  p50 {kvc_p50:.0f}  p90 {np.quantile(kvc_x, 0.9):.0f}")
+
+    print(f"\nDP 4w:   total_input={sum(r['input_length'] for r in dp)/1e6:.1f}M tokens")
+    print(f"         total_cached={sum(r.get('cached_tokens',0) for r in dp)/1e6:.1f}M tokens")
+    print(f"         hit rate {sum(r.get('cached_tokens',0) for r in dp)/sum(r['input_length'] for r in dp)*100:.2f}%")
+    print(f"         mean uncached {np.mean(dp_x):.0f}  p50 {dp_p50:.0f}  p90 {np.quantile(dp_x, 0.9):.0f}")
+
+    print(f"\nMid-turn region (8-25): KVC {drift_kvc*100:.2f}%  DP {drift_dp*100:.2f}%  (gap {(drift_kvc-drift_dp)*100:+.2f}pp)")
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/plot_gpu_utilization.py

@@ -0,0 +1,234 @@
+#!/usr/bin/env python3
+"""Per-GPU utilization breakdown: KVC 1P3D v2 vs 4-way DP CA.
+
+Generates docs/figures/gpu_utilization.png — two-panel:
+  left:  per-GPU request count
+  right: per-GPU compute work (uncached prefill tokens + decode tokens, stacked)
+
+The point of the figure is to push back on the naïve reading
+"KVC's prefill GPU is idle 90% of the time, so KVC is using fewer GPUs."
+
+By request count, the prefill GPU is indeed touched by only ~8% of requests.
+By compute work, the prefill GPU bears comparable per-GPU load to each
+decode GPU — it is a low-frequency, high-cost safety net for cache misses,
+not idle capacity.
+
+Work attribution:
+  KVC direct-to-D path: prefill happens locally on the assigned D worker
+                        (append-prefill of `uncached_tokens` tokens).
+  KVC seed/reseed/fallback path: prefill happens on prefill-0
+                        (full uncached_tokens), decode on assigned D.
+  DP: all work on assigned direct-N worker.
+
+Aborted / errored requests are excluded.
+"""
+
+from __future__ import annotations
+
+import json
+from collections import defaultdict
+from pathlib import Path
+
+import matplotlib.pyplot as plt
+import numpy as np
+
+ROOT = Path(__file__).resolve().parents[2]
+KVC = ROOT / "outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl"
+DP = ROOT / "outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl"
+OUT = ROOT / "docs/figures/gpu_utilization.png"
+
+
+def load(p: Path) -> list[dict]:
+    return [json.loads(line) for line in p.open()]
+
+
+def is_failed(r: dict) -> bool:
+    if r.get("error"):
+        return True
+    fr = r.get("finish_reason")
+    if fr and ("abort" in str(fr).lower() or "badrequest" in str(fr).lower()):
+        return True
+    return False
+
+
+def uncached(r: dict) -> int:
+    return max(0, r["input_length"] - r.get("cached_tokens", 0))
+
+
+def out_tokens(r: dict) -> int:
+    return r.get("actual_output_tokens") or r.get("output_length") or 0
+
+
+def main() -> None:
+    kvc = [r for r in load(KVC) if not is_failed(r)]
+    dp = [r for r in load(DP) if not is_failed(r)]
+
+    # ------------------------------------------------------------------
+    # KVC per-GPU attribution
+    # ------------------------------------------------------------------
+    kvc_req_count = defaultdict(int)
+    kvc_prefill_tokens = defaultdict(int)   # uncached prefill compute
+    kvc_decode_tokens = defaultdict(int)
+
+    for r in kvc:
+        d = r["assigned_decode_node"]            # decode-0/1/2
+        p = r["assigned_prefill_node"]            # prefill-0
+        mode = r.get("execution_mode", "")
+        if mode == "kvcache-direct-to-d-session":
+            # P is bypassed entirely; D does the append-prefill + decode
+            kvc_req_count[d] += 1
+            kvc_prefill_tokens[d] += uncached(r)
+            kvc_decode_tokens[d] += out_tokens(r)
+        else:
+            # P does the full prefill; D handles decode
+            kvc_req_count[p] += 1
+            kvc_req_count[d] += 1   # decode side still counts
+            kvc_prefill_tokens[p] += uncached(r)
+            kvc_decode_tokens[d] += out_tokens(r)
+
+    # ------------------------------------------------------------------
+    # DP per-GPU attribution (fused P+D on every worker)
+    # ------------------------------------------------------------------
+    dp_req_count = defaultdict(int)
+    dp_prefill_tokens = defaultdict(int)
+    dp_decode_tokens = defaultdict(int)
+
+    for r in dp:
+        w = r["assigned_decode_node"]  # direct-0..3
+        dp_req_count[w] += 1
+        dp_prefill_tokens[w] += uncached(r)
+        dp_decode_tokens[w] += out_tokens(r)
+
+    # ------------------------------------------------------------------
+    # Build ordered GPU list, KVC then DP
+    # ------------------------------------------------------------------
+    kvc_gpus = ["prefill-0", "decode-0", "decode-1", "decode-2"]
+    dp_gpus = ["direct-0", "direct-1", "direct-2", "direct-3"]
+    all_gpus = kvc_gpus + dp_gpus
+
+    def get(d, k):
+        return d.get(k, 0)
+
+    counts = [get(kvc_req_count, g) for g in kvc_gpus] + \
+             [get(dp_req_count, g) for g in dp_gpus]
+    prefill_tk = [get(kvc_prefill_tokens, g) for g in kvc_gpus] + \
+                 [get(dp_prefill_tokens, g) for g in dp_gpus]
+    decode_tk = [get(kvc_decode_tokens, g) for g in kvc_gpus] + \
+                [get(dp_decode_tokens, g) for g in dp_gpus]
+
+    # Display labels: P/D role + worker id
+    labels = [
+        "KVC P\nprefill-0",
+        "KVC D\ndecode-0",
+        "KVC D\ndecode-1",
+        "KVC D\ndecode-2",
+        "DP P+D\ndirect-0",
+        "DP P+D\ndirect-1",
+        "DP P+D\ndirect-2",
+        "DP P+D\ndirect-3",
+    ]
+    kvc_mask = [True, True, True, True, False, False, False, False]
+
+    KVC_P_COLOR = "#E89D44"     # orange — P GPU stands out
+    KVC_D_COLOR = "#1F77B4"     # blue
+    DP_COLOR    = "#D62728"     # red
+
+    bar_colors = [KVC_P_COLOR, KVC_D_COLOR, KVC_D_COLOR, KVC_D_COLOR,
+                  DP_COLOR, DP_COLOR, DP_COLOR, DP_COLOR]
+
+    fig, axes = plt.subplots(1, 2, figsize=(15, 6.5))
+    x = np.arange(len(all_gpus))
+
+    # -- Left: per-GPU request count ----------------------------------
+    ax = axes[0]
+    bars = ax.bar(x, counts, color=bar_colors, edgecolor="black", linewidth=0.6)
+    for xi, c in zip(x, counts):
+        ax.text(xi, c + max(counts) * 0.015, f"{c:,}",
+                ha="center", va="bottom", fontsize=9.5)
+    ax.set_xticks(x)
+    ax.set_xticklabels(labels, fontsize=9.5)
+    ax.set_ylabel("Number of requests touching this GPU", fontsize=11)
+    ax.set_title("Per-GPU request count\n(naïve view: P seems idle)", fontsize=12, pad=10)
+    ax.grid(axis="y", linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+
+    # Annotate: KVC P GPU is "low frequency"
+    p_idx = 0
+    p_pct = counts[p_idx] / sum(counts[:4]) * 100  # vs KVC total
+    ax.annotate(
+        f"P GPU only sees\n"
+        f"{counts[p_idx]:,} requests\n"
+        f"({counts[p_idx]/len(kvc)*100:.1f}% of total)",
+        xy=(p_idx, counts[p_idx]),
+        xytext=(p_idx + 0.6, max(counts) * 0.55),
+        fontsize=9, color=KVC_P_COLOR, fontweight="bold",
+        arrowprops=dict(arrowstyle="->", color=KVC_P_COLOR, lw=1.0),
+    )
+
+    # -- Right: per-GPU compute work (stacked prefill + decode) -------
+    ax = axes[1]
+    prefill_M = [t / 1e6 for t in prefill_tk]
+    decode_M = [t / 1e6 for t in decode_tk]
+    total_M = [p + d for p, d in zip(prefill_M, decode_M)]
+
+    bars_p = ax.bar(x, prefill_M, color=[c for c in bar_colors],
+                    edgecolor="black", linewidth=0.6, label="Uncached prefill tokens",
+                    alpha=0.95)
+    bars_d = ax.bar(x, decode_M, bottom=prefill_M, color=[c for c in bar_colors],
+                    edgecolor="black", linewidth=0.6, hatch="///",
+                    label="Decode tokens", alpha=0.55)
+
+    for xi, t in zip(x, total_M):
+        ax.text(xi, t + max(total_M) * 0.015, f"{t:.2f}M",
+                ha="center", va="bottom", fontsize=9.5)
+
+    ax.set_xticks(x)
+    ax.set_xticklabels(labels, fontsize=9.5)
+    ax.set_ylabel("Compute tokens (millions)", fontsize=11)
+    ax.set_title("Per-GPU compute work\n(work view: P is comparable to each D)",
+                 fontsize=12, pad=10)
+    ax.grid(axis="y", linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+    ax.legend(loc="upper left", fontsize=10, framealpha=0.95)
+
+    # Annotate: KVC P GPU does similar work to each D
+    ax.annotate(
+        f"P GPU does {total_M[p_idx]:.2f}M tokens of\n"
+        f"prefill — comparable per-GPU\n"
+        f"load to each KVC D worker",
+        xy=(p_idx, total_M[p_idx]),
+        xytext=(p_idx + 0.6, max(total_M) * 0.62),
+        fontsize=9, color=KVC_P_COLOR, fontweight="bold",
+        arrowprops=dict(arrowstyle="->", color=KVC_P_COLOR, lw=1.0),
+    )
+
+    # Separator + group labels
+    for ax in axes:
+        ax.axvline(3.5, color="gray", linestyle="--", linewidth=1.0, alpha=0.5)
+        ymin, ymax = ax.get_ylim()
+        ax.text(1.5, ymax * 1.05, "KVC 1P3D", ha="center", fontsize=11,
+                fontweight="bold", color="#444")
+        ax.text(5.5, ymax * 1.05, "DP 4-way CA", ha="center", fontsize=11,
+                fontweight="bold", color="#444")
+
+    fig.suptitle(
+        "Per-GPU utilization: \"is KVC's prefill GPU wasted?\"\n"
+        "Left view says yes (only 8% of requests); right view says no (comparable work to each D).",
+        fontsize=13, y=1.02,
+    )
+    plt.tight_layout()
+    plt.savefig(OUT, dpi=150, bbox_inches="tight")
+    print(f"wrote {OUT}")
+    plt.close(fig)
+
+    # ------------------------------------------------------------------
+    # Print numbers for doc reference
+    # ------------------------------------------------------------------
+    print("\n=== Per-GPU numbers ===")
+    print(f"{'GPU':<22}  {'requests':>10}  {'prefill(M)':>12}  {'decode(M)':>12}  {'total(M)':>10}")
+    for lbl, n, pM, dM in zip(labels, counts, prefill_M, decode_M):
+        print(f"  {lbl.replace(chr(10), ' '):<20}  {n:>10}  {pM:>12.3f}  {dM:>12.3f}  {pM+dM:>10.3f}")
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/plot_ttft_pdf.py

@@ -0,0 +1,199 @@
+#!/usr/bin/env python3
+"""Generate TTFT probability density curves: KVC 1P3D v2 vs 4-way DP CA.
+
+Inputs:
+  outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl
+  outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl
+
+Outputs:
+  docs/figures/ttft_pdf_comparison.png  -- two-panel figure:
+      left panel: linear x in [0, 1.0]s zoomed on the body
+      right panel: log x covering full range (0.01 -- 10 s)
+  Each KDE curve uses scipy.stats.gaussian_kde with Scott's rule bandwidth.
+  Aborted requests are excluded (same filter as metrics.py:_is_failed_request).
+"""
+
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import matplotlib.pyplot as plt
+import numpy as np
+from scipy.stats import gaussian_kde
+
+ROOT = Path(__file__).resolve().parents[2]
+KVC = ROOT / "outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl"
+DP = ROOT / "outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl"
+OUT = ROOT / "docs/figures/ttft_pdf_comparison.png"
+
+
+def load(p: Path) -> list[dict]:
+    return [json.loads(line) for line in p.open()]
+
+
+def is_failed(r: dict) -> bool:
+    if r.get("error"):
+        return True
+    fr = r.get("finish_reason")
+    if fr and ("abort" in str(fr).lower() or "badrequest" in str(fr).lower()):
+        return True
+    return False
+
+
+def pct(vals: np.ndarray, q: float) -> float:
+    return float(np.quantile(vals, q))
+
+
+def main() -> None:
+    kvc = [r for r in load(KVC) if not is_failed(r)]
+    dp = [r for r in load(DP) if not is_failed(r)]
+
+    kvc_ttft = np.array([r["ttft_s"] for r in kvc if r.get("ttft_s") is not None])
+    dp_ttft = np.array([r["ttft_s"] for r in dp if r.get("ttft_s") is not None])
+
+    # Trim absurdly small zeros (rare measurement artifacts) so log KDE behaves.
+    kvc_ttft = kvc_ttft[kvc_ttft > 1e-4]
+    dp_ttft = dp_ttft[dp_ttft > 1e-4]
+
+    KVC_COLOR = "#1F77B4"  # blue
+    DP_COLOR = "#D62728"   # red
+
+    fig, axes = plt.subplots(1, 2, figsize=(16, 6.5))
+
+    # ------------------------------------------------------------------
+    # Left panel: linear x ∈ [0, 0.6]s -- body of the distribution
+    # ------------------------------------------------------------------
+    ax = axes[0]
+    x_body = np.linspace(0.0, 0.6, 600)
+
+    # KDE on linear ttft values, clipped to body
+    kde_kvc_lin = gaussian_kde(kvc_ttft, bw_method=0.15)
+    kde_dp_lin = gaussian_kde(dp_ttft, bw_method=0.15)
+
+    ax.plot(x_body, kde_kvc_lin(x_body),
+            color=KVC_COLOR, lw=2.5, label=f"KVC 1P3D v2  (n={len(kvc_ttft)})")
+    ax.fill_between(x_body, kde_kvc_lin(x_body), alpha=0.20, color=KVC_COLOR)
+    ax.plot(x_body, kde_dp_lin(x_body),
+            color=DP_COLOR, lw=2.5, label=f"4-way DP CA  (n={len(dp_ttft)})")
+    ax.fill_between(x_body, kde_dp_lin(x_body), alpha=0.20, color=DP_COLOR)
+
+    # Vertical lines for p50, p90
+    for q, ls in [(0.50, "-"), (0.90, "--")]:
+        ax.axvline(pct(kvc_ttft, q), color=KVC_COLOR, ls=ls, alpha=0.55, lw=1.1)
+        ax.axvline(pct(dp_ttft, q), color=DP_COLOR, ls=ls, alpha=0.55, lw=1.1)
+    ymax = ax.get_ylim()[1]
+    ax.text(pct(kvc_ttft, 0.50), ymax * 0.97,
+            f"KVC p50\n{pct(kvc_ttft, 0.50)*1000:.0f}ms",
+            color=KVC_COLOR, fontsize=9, va="top", ha="left",
+            bbox=dict(facecolor="white", edgecolor="none", alpha=0.7, pad=2))
+    ax.text(pct(dp_ttft, 0.50), ymax * 0.50,
+            f"DP p50\n{pct(dp_ttft, 0.50)*1000:.0f}ms",
+            color=DP_COLOR, fontsize=9, va="top", ha="left",
+            bbox=dict(facecolor="white", edgecolor="none", alpha=0.7, pad=2))
+    ax.text(pct(kvc_ttft, 0.90), ymax * 0.30,
+            f"KVC p90\n{pct(kvc_ttft, 0.90)*1000:.0f}ms",
+            color=KVC_COLOR, fontsize=9, va="top", ha="left",
+            bbox=dict(facecolor="white", edgecolor="none", alpha=0.7, pad=2))
+    ax.text(pct(dp_ttft, 0.90), ymax * 0.18,
+            f"DP p90\n{pct(dp_ttft, 0.90)*1000:.0f}ms",
+            color=DP_COLOR, fontsize=9, va="top", ha="left",
+            bbox=dict(facecolor="white", edgecolor="none", alpha=0.7, pad=2))
+
+    ax.set_xlim(0, 0.6)
+    ax.set_xlabel("TTFT (seconds, linear)", fontsize=11)
+    ax.set_ylabel("Probability density", fontsize=11)
+    ax.set_title("Body of distribution  (TTFT ≤ 0.6 s)", fontsize=12, pad=10)
+    ax.legend(loc="upper right", fontsize=10, framealpha=0.95)
+    ax.grid(True, linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+
+    # ------------------------------------------------------------------
+    # Right panel: log x ∈ [0.01, 10]s -- full range incl. tail
+    # PDF on log-x: we plot density vs log10(t) so the curve integrates
+    # to 1 over log space (standard "log-density" presentation).
+    # ------------------------------------------------------------------
+    ax = axes[1]
+    # KDE on log10(ttft) so the resulting curve integrates to 1 over log10 t
+    kde_kvc_log = gaussian_kde(np.log10(kvc_ttft), bw_method="scott")
+    kde_dp_log = gaussian_kde(np.log10(dp_ttft), bw_method="scott")
+    log_x = np.linspace(np.log10(0.01), np.log10(10.0), 600)
+    x_full = 10 ** log_x
+
+    y_kvc = kde_kvc_log(log_x)
+    y_dp = kde_dp_log(log_x)
+
+    ax.plot(x_full, y_kvc, color=KVC_COLOR, lw=2.5, label=f"KVC 1P3D v2  (n={len(kvc_ttft)})")
+    ax.fill_between(x_full, y_kvc, alpha=0.20, color=KVC_COLOR)
+    ax.plot(x_full, y_dp, color=DP_COLOR, lw=2.5, label=f"4-way DP CA  (n={len(dp_ttft)})")
+    ax.fill_between(x_full, y_dp, alpha=0.20, color=DP_COLOR)
+
+    ax.set_xscale("log")
+    ax.set_xlim(0.01, 10.0)
+
+    # Percentile markers
+    quartile_styles = [(0.50, "-", "p50"), (0.90, "--", "p90"), (0.99, ":", "p99")]
+    for q, ls, name in quartile_styles:
+        ax.axvline(pct(kvc_ttft, q), color=KVC_COLOR, ls=ls, alpha=0.55, lw=1.1)
+        ax.axvline(pct(dp_ttft, q), color=DP_COLOR, ls=ls, alpha=0.55, lw=1.1)
+
+    # Annotate p99 specifically since this is the key reviewer-targeted callout
+    ymax = max(y_kvc.max(), y_dp.max())
+    kvc_p99 = pct(kvc_ttft, 0.99)
+    dp_p99 = pct(dp_ttft, 0.99)
+    ax.annotate(f"KVC p99 = {kvc_p99:.2f}s\n(slow-path reseed tail)",
+                xy=(kvc_p99, kde_kvc_log(np.log10(kvc_p99))[0]),
+                xytext=(2.0, ymax * 0.65),
+                fontsize=10, color=KVC_COLOR, fontweight="bold",
+                arrowprops=dict(arrowstyle="->", color=KVC_COLOR, lw=1.0))
+    ax.annotate(f"DP p99 = {dp_p99*1000:.0f}ms",
+                xy=(dp_p99, kde_dp_log(np.log10(dp_p99))[0]),
+                xytext=(0.025, ymax * 0.80),
+                fontsize=10, color=DP_COLOR, fontweight="bold",
+                arrowprops=dict(arrowstyle="->", color=DP_COLOR, lw=1.0))
+    # Highlight the KVC bimodal structure
+    ax.annotate("KVC fast path\n(direct-to-D, 91.6%)",
+                xy=(0.05, y_kvc[np.argmin(np.abs(x_full - 0.05))]),
+                xytext=(0.012, ymax * 0.45),
+                fontsize=9, color=KVC_COLOR, style="italic",
+                arrowprops=dict(arrowstyle="->", color=KVC_COLOR, lw=0.7, alpha=0.6))
+    ax.annotate("KVC slow path\n(reseed, ~3.4%)",
+                xy=(2.5, y_kvc[np.argmin(np.abs(x_full - 2.5))]),
+                xytext=(3.0, ymax * 0.30),
+                fontsize=9, color=KVC_COLOR, style="italic",
+                arrowprops=dict(arrowstyle="->", color=KVC_COLOR, lw=0.7, alpha=0.6))
+
+    # Custom tick labels in seconds (instead of 10^-2, 10^-1, 10^0, 10^1)
+    ax.set_xticks([0.01, 0.05, 0.1, 0.5, 1.0, 5.0, 10.0])
+    ax.set_xticklabels(["10ms", "50ms", "100ms", "500ms", "1s", "5s", "10s"])
+
+    ax.set_xlabel("TTFT (log scale)", fontsize=11)
+    ax.set_ylabel("Density  (per log₁₀ s)", fontsize=11)
+    ax.set_title("Full range  (TTFT 10 ms – 10 s, log x)", fontsize=12, pad=10)
+    ax.legend(loc="upper left", fontsize=10, framealpha=0.95)
+    ax.grid(True, which="both", linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+
+    fig.suptitle(
+        "TTFT probability density: KVC 1P3D v2 vs 4-way DP CA\n"
+        "SWE-Bench 50sess trace · ts=1 · 4× H100 80GB · aborted/error requests excluded",
+        fontsize=13, y=1.02,
+    )
+    plt.tight_layout()
+    plt.savefig(OUT, dpi=150, bbox_inches="tight")
+    print(f"wrote {OUT}")
+    plt.close(fig)
+
+    # ------------------------------------------------------------------
+    # Print summary stats for doc cross-reference
+    # ------------------------------------------------------------------
+    print(f"\n=== TTFT distribution summary ===")
+    for name, arr in [("KVC v2", kvc_ttft), ("DP 4w", dp_ttft)]:
+        print(f"  {name}  (n={len(arr)})")
+        print(f"    min={arr.min()*1000:.1f}ms  p10={pct(arr,0.10)*1000:.1f}ms  "
+              f"p50={pct(arr,0.50)*1000:.1f}ms  p90={pct(arr,0.90)*1000:.1f}ms  "
+              f"p99={pct(arr,0.99)*1000:.1f}ms  max={arr.max()*1000:.1f}ms")
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/plot_v2_path_breakdown.py

@@ -0,0 +1,223 @@
+#!/usr/bin/env python3
+"""Generate the two figures referenced by docs/V2_DEEP_ANALYSIS_ZH.md §3.1 and §3.2.
+
+Inputs:
+  outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl
+  outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl
+
+Outputs:
+  docs/figures/v2_execution_mode_distribution.png   (for §3.1)
+  docs/figures/v2_path_level_latency.png            (for §3.2)
+"""
+
+from __future__ import annotations
+
+import json
+import statistics
+from collections import Counter, defaultdict
+from pathlib import Path
+
+import matplotlib.pyplot as plt
+import numpy as np
+
+ROOT = Path(__file__).resolve().parents[2]
+KVC = ROOT / "outputs/qwen3-30b-tp1-ts1-migration-v2/kvc_1p3d_migration_v2_run1_metrics.jsonl"
+DP = ROOT / "outputs/qwen3-30b-tp1-ts1-validation/dp4_metrics.jsonl"
+OUT = ROOT / "docs/figures"
+OUT.mkdir(parents=True, exist_ok=True)
+
+
+def load(p: Path) -> list[dict]:
+    return [json.loads(line) for line in p.open()]
+
+
+def is_failed(r: dict) -> bool:
+    if r.get("error"):
+        return True
+    fr = r.get("finish_reason")
+    if fr and ("abort" in str(fr).lower() or "badrequest" in str(fr).lower()):
+        return True
+    return False
+
+
+def pct(vals: list[float], q: float) -> float:
+    s = sorted(vals)
+    if not s:
+        return float("nan")
+    return s[max(0, min(len(s) - 1, int(len(s) * q)))]
+
+
+def main() -> None:
+    kvc = load(KVC)
+    dp = load(DP)
+
+    kvc_ok = [r for r in kvc if not is_failed(r)]
+    dp_ok = [r for r in dp if not is_failed(r)]
+
+    # ------------------------------------------------------------------
+    # Figure 1: §3.1 execution_mode distribution (horizontal bar)
+    # Use ALL rows (incl. failures) so percentages match the doc's 91.6%
+    # ------------------------------------------------------------------
+    mode_counts = Counter(r["execution_mode"] for r in kvc)
+    total_kvc = len(kvc)
+
+    short_label = {
+        "kvcache-direct-to-d-session": "direct-to-D-session  (fast path)",
+        "pd-router-d-session-reseed": "d-session-reseed  (mooncake reseed)",
+        "pd-router-fallback-session-not-resident-session-cap":
+            "fallback: session-not-resident + session-cap",
+        "pd-router-fallback-session-not-resident-seed-filter-early-turn":
+            "fallback: session-not-resident + seed-filter",
+        "pd-router-turn1-seed": "turn1-seed  (first turn of each session)",
+        "pd-router-fallback-no-d-capacity": "fallback: no-d-capacity",
+        "pd-router-fallback-real-large-append-session-cap":
+            "fallback: real-large-append",
+        "pd-router-fallback-policy-no-bypass-session-cap":
+            "fallback: policy-no-bypass",
+        "pd-router-d-session-reseed-after-eviction":
+            "d-session-reseed-after-eviction",
+        "kvcache-centric": "kvcache-centric (admit-but-then-error)",
+    }
+    sorted_modes = mode_counts.most_common()
+    labels = [short_label.get(m, m) for m, _ in sorted_modes]
+    counts = [c for _, c in sorted_modes]
+    pcts = [c / total_kvc * 100 for c in counts]
+
+    is_fast = ["direct-to-D" in lbl for lbl in labels]
+    colors = ["#2C8C2C" if f else "#D62728" for f in is_fast]
+
+    fig, ax = plt.subplots(figsize=(11, 5.5))
+    y = np.arange(len(labels))[::-1]
+    ax.barh(y, counts, color=colors, edgecolor="black", linewidth=0.5)
+    ax.set_yticks(y)
+    ax.set_yticklabels(labels, fontsize=10)
+    ax.set_xscale("log")
+    ax.set_xlabel("Request count (log scale)", fontsize=11)
+    ax.set_xlim(left=1)
+
+    # Annotate count + percentage at end of each bar
+    for yi, (c, p) in zip(y, zip(counts, pcts)):
+        ax.text(c * 1.05, yi, f"{c}  ({p:.1f}%)",
+                va="center", fontsize=9.5)
+
+    ax.set_title(
+        f"KVC v2 execution_mode distribution  (n = {total_kvc} total requests)\n"
+        "green = fast path (direct-to-D), red = slow / fallback / failure paths",
+        fontsize=12, pad=12,
+    )
+    ax.grid(axis="x", linestyle=":", alpha=0.4)
+    ax.set_axisbelow(True)
+    plt.tight_layout()
+    out1 = OUT / "v2_execution_mode_distribution.png"
+    plt.savefig(out1, dpi=150)
+    print(f"wrote {out1}")
+    plt.close(fig)
+
+    # ------------------------------------------------------------------
+    # Figure 2: §3.2 path-level latency (grouped bars, log y)
+    # ------------------------------------------------------------------
+
+    # Group KVC paths semantically
+    def kvc_group(mode: str) -> str:
+        if mode == "kvcache-direct-to-d-session":
+            return "KVC direct-to-D\n(fast path, 91.6%)"
+        if "reseed" in mode:
+            return "KVC reseed\n(slow path, 3.4%)"
+        if "no-d-capacity" in mode:
+            return "KVC no-d-capacity\n(fallback, 0.7%)"
+        if "session-not-resident" in mode:
+            return "KVC session-not-resident\n(misc, 2.3%)"
+        return "KVC other\n(<2%)"
+
+    groups = defaultdict(list)
+    for r in kvc_ok:
+        groups[kvc_group(r["execution_mode"])].append(r)
+
+    # Order paths by intuitive progression (fast → slow)
+    ordered_paths = [
+        "KVC direct-to-D\n(fast path, 91.6%)",
+        "KVC session-not-resident\n(misc, 2.3%)",
+        "KVC reseed\n(slow path, 3.4%)",
+        "KVC no-d-capacity\n(fallback, 0.7%)",
+    ]
+    # Filter to only ones present
+    ordered_paths = [p for p in ordered_paths if p in groups]
+    ordered_paths.append("DP dp-colo-router\n(100%)")
+
+    def stats(rows: list[dict]) -> dict[str, float]:
+        ttfts = [r["ttft_s"] for r in rows if r.get("ttft_s") is not None]
+        lats = [r["latency_s"] for r in rows if r.get("latency_s") is not None]
+        return {
+            "n": len(rows),
+            "ttft_p50": pct(ttfts, 0.50),
+            "ttft_p99": pct(ttfts, 0.99),
+            "lat_p50": pct(lats, 0.50),
+        }
+
+    path_stats = {p: stats(groups[p]) for p in ordered_paths if "DP" not in p}
+    path_stats["DP dp-colo-router\n(100%)"] = stats(dp_ok)
+
+    metrics = [("TTFT p50", "ttft_p50"), ("TTFT p99", "ttft_p99"), ("Latency p50", "lat_p50")]
+    bar_w = 0.25
+    fig, ax = plt.subplots(figsize=(12, 6))
+    x = np.arange(len(ordered_paths))
+
+    colors_metric = ["#1F77B4", "#FF7F0E", "#9467BD"]
+    for i, (label, key) in enumerate(metrics):
+        vals = [path_stats[p][key] for p in ordered_paths]
+        bars = ax.bar(x + (i - 1) * bar_w, vals, bar_w, label=label,
+                      color=colors_metric[i], edgecolor="black", linewidth=0.4)
+        for xi, v in zip(x + (i - 1) * bar_w, vals):
+            if v > 0 and v == v:  # not nan
+                fmt = f"{v*1000:.0f}ms" if v < 1 else f"{v:.2f}s"
+                ax.text(xi, v * 1.10, fmt,
+                        ha="center", va="bottom", fontsize=8.5, rotation=0)
+
+    ax.set_yscale("log")
+    ax.set_xticks(x)
+    ax.set_xticklabels(ordered_paths, fontsize=9.5)
+    ax.set_ylabel("Latency (seconds, log scale)", fontsize=11)
+    ax.set_title(
+        "Path-level latency: KVC v2 paths vs DP single-path baseline\n"
+        "log y-axis · same SWE-Bench 50sess trace · ts=1 · 4× H100 80GB",
+        fontsize=12, pad=12,
+    )
+    ax.legend(loc="upper left", fontsize=10, framealpha=0.95)
+    ax.grid(axis="y", linestyle=":", alpha=0.4, which="both")
+    ax.set_axisbelow(True)
+
+    # Annotate sample counts under each path label
+    ymin = ax.get_ylim()[0]
+    for xi, p in zip(x, ordered_paths):
+        n = path_stats[p]["n"]
+        ax.text(xi, ymin * 0.5, f"n={n}", ha="center", va="top",
+                fontsize=8.5, color="#555")
+
+    plt.tight_layout()
+    out2 = OUT / "v2_path_level_latency.png"
+    plt.savefig(out2, dpi=150)
+    print(f"wrote {out2}")
+    plt.close(fig)
+
+    # ------------------------------------------------------------------
+    # Print numeric values used (for doc reference)
+    # ------------------------------------------------------------------
+    print("\n=== Numeric values plotted ===")
+    print("\nExecution mode counts (KVC v2):")
+    for label, c, p in zip(labels, counts, pcts):
+        print(f"  {c:>5}  ({p:>5.2f}%)  {label}")
+
+    print("\nPath-level latency:")
+    for p in ordered_paths:
+        s = path_stats[p]
+        nl = " | ".join([
+            f"n={s['n']}",
+            f"TTFT p50={s['ttft_p50']*1000:.1f}ms",
+            f"TTFT p99={s['ttft_p99']*1000:.1f}ms",
+            f"Lat p50={s['lat_p50']:.3f}s",
+        ])
+        print(f"  {p.replace(chr(10), ' '):<55}  {nl}")
+
+
+if __name__ == "__main__":
+    main()

scripts/analysis/recompute_summary.py

@@ -0,0 +1,105 @@
+#!/usr/bin/env python3
+"""Re-derive summary.json from existing metrics.jsonl using the fixed metrics.py.
+
+Bug fixed: requests aborted by SGLang (e.g. input > max-input-len returns
+a fast 400 with latency_s ~ 0.08s) were previously counted in latency_stats
+as if successful, deflating mean/p50/p90. The fixed metrics.py excludes
+all failed requests (errors or aborts) from latency/ttft/tpot stats and
+exposes abort_count / failure_count.
+
+Usage:
+    python3 scripts/analysis/recompute_summary.py path/to/metrics.jsonl ...
+    python3 scripts/analysis/recompute_summary.py --diff path/to/metrics.jsonl path/to/old_summary.json
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+import sys
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[2] / "src"))
+
+from agentic_pd_hybrid.metrics import RequestMetrics, write_summary_json
+
+
+def load_rows(metrics_path: Path) -> list[RequestMetrics]:
+    rows = []
+    field_names = {f for f in RequestMetrics.__dataclass_fields__}
+    with metrics_path.open() as handle:
+        for line in handle:
+            line = line.strip()
+            if not line:
+                continue
+            raw = json.loads(line)
+            kwargs = {k: raw.get(k) for k in field_names}
+            rows.append(RequestMetrics(**kwargs))
+    return rows
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("metrics_paths", nargs="+", type=Path)
+    parser.add_argument(
+        "--out",
+        type=Path,
+        default=None,
+        help="output summary path (default: alongside metrics with .recomputed_summary.json)",
+    )
+    parser.add_argument(
+        "--diff",
+        action="store_true",
+        help="print before/after diff against the old <metrics>.summary.json",
+    )
+    args = parser.parse_args()
+
+    for metrics_path in args.metrics_paths:
+        rows = load_rows(metrics_path)
+        out_path = args.out or metrics_path.with_suffix(".recomputed_summary.json")
+        write_summary_json(
+            out_path,
+            rows,
+            trace_path=metrics_path,
+            router_url=None,
+        )
+        new = json.load(out_path.open())
+        print(f"\n=== {metrics_path} ===")
+        print(f"  written: {out_path}")
+        print(f"  total rows:     {new['request_count']}")
+        print(f"  error_count:    {new['error_count']}")
+        print(f"  abort_count:    {new.get('abort_count', '?')}")
+        print(f"  failure_count:  {new.get('failure_count', '?')}")
+        ls = new.get("latency_stats_s", {}) or {}
+        ts = new.get("ttft_stats_s", {}) or {}
+        print(f"  lat:  n={ls.get('count')} mean={ls.get('mean'):.4f} p50={ls.get('p50'):.4f} p90={ls.get('p90'):.4f} p99={ls.get('p99'):.4f}")
+        print(f"  ttft: n={ts.get('count')} mean={ts.get('mean'):.4f} p50={ts.get('p50'):.4f} p90={ts.get('p90'):.4f} p99={ts.get('p99'):.4f}")
+
+        if args.diff:
+            # find old summary (sibling file)
+            candidates = [
+                metrics_path.parent / f"{metrics_path.stem}.summary.json",
+                metrics_path.with_suffix(".summary.json"),
+            ]
+            old_path = next((p for p in candidates if p.exists()), None)
+            if old_path:
+                old = json.load(old_path.open())
+                print(f"  vs old {old_path}:")
+                old_ls = old.get("latency_stats_s", {}) or {}
+                old_ts = old.get("ttft_stats_s", {}) or {}
+                for k in ("count", "mean", "p50", "p90", "p99"):
+                    o = old_ls.get(k)
+                    n = ls.get(k)
+                    if o is not None and n is not None:
+                        delta = n - o
+                        print(f"    lat.{k}:  {o:.4f} -> {n:.4f}  ({delta:+.4f})")
+                for k in ("count", "mean", "p50", "p90", "p99"):
+                    o = old_ts.get(k)
+                    n = ts.get(k)
+                    if o is not None and n is not None:
+                        delta = n - o
+                        print(f"    ttft.{k}: {o:.4f} -> {n:.4f}  ({delta:+.4f})")
+
+
+if __name__ == "__main__":
+    main()

scripts/convert_audit_to_trace.py

@@ -0,0 +1,110 @@
+#!/usr/bin/env python3
+"""Convert sibench audit.jsonl to agentic-pd-hybrid trace format.
+
+Source format (sibench audit.jsonl):
+  {"instance_id": "...", "ts": float, "messages": [...],
+   "audit": {"prompt_tokens": int, "completion_tokens": int, ...}}
+
+Target format (agentic-pd-hybrid trace JSONL):
+  {"chat_id": int, "parent_chat_id": int, "timestamp": float,
+   "turn": int, "input_length": int, "output_length": int,
+   "type": str, "hash_ids": [int, ...]}
+"""
+
+import json
+import sys
+from collections import defaultdict
+from pathlib import Path
+
+BLOCK_TOKEN_BUDGET = 24  # tokens per block, matching trace.py default
+
+
+def convert(src: Path, dst: Path) -> None:
+    # Group lines by instance_id, preserving order within each instance
+    instances: dict[str, list[dict]] = defaultdict(list)
+    with src.open() as f:
+        for line in f:
+            line = line.strip()
+            if not line:
+                continue
+            rec = json.loads(line)
+            instances[rec["instance_id"]].append(rec)
+
+    # Sort each instance's turns by timestamp
+    for iid in instances:
+        instances[iid].sort(key=lambda r: r["ts"])
+
+    # Assign stable chat_id bases: each instance gets a block of IDs
+    # Max turns across all instances determines the spacing
+    max_turns = max(len(turns) for turns in instances.values())
+    spacing = max_turns + 10  # extra headroom
+
+    total_written = 0
+    with dst.open("w") as out:
+        for inst_idx, (iid, turns) in enumerate(instances.items()):
+            base_chat_id = (inst_idx + 1) * spacing  # start from spacing to avoid 0
+            # Track cumulative hash_ids for prefix cache simulation
+            cumulative_hash_ids: list[int] = []
+            global_block_counter = inst_idx * 100_000  # unique block namespace per instance
+
+            for turn_idx, rec in enumerate(turns):
+                audit = rec.get("audit", {})
+                input_length = audit.get("prompt_tokens", 0)
+                output_length = audit.get("completion_tokens", 0)
+
+                if input_length <= 0:
+                    # Fallback: estimate from message content
+                    total_chars = sum(len(m.get("content", "")) for m in rec.get("messages", []))
+                    input_length = max(1, total_chars // 4)
+                if output_length <= 0:
+                    output_length = 128  # reasonable default
+
+                chat_id = base_chat_id + turn_idx
+                if turn_idx == 0:
+                    parent_chat_id = -1
+                else:
+                    parent_chat_id = base_chat_id + turn_idx - 1
+
+                # Build hash_ids: for turn 0, generate blocks for full input
+                # For turn N>0, keep previous blocks and add new ones for the delta
+                if turn_idx == 0:
+                    num_blocks = input_length // BLOCK_TOKEN_BUDGET
+                    cumulative_hash_ids = list(
+                        range(global_block_counter, global_block_counter + num_blocks)
+                    )
+                    global_block_counter += num_blocks
+                else:
+                    # The new input is the full prompt (cumulative), so the delta
+                    # is the new tokens beyond what was in the previous turn's prompt
+                    prev_input = audit.get("prompt_tokens", 0)
+                    prev_rec_audit = turns[turn_idx - 1].get("audit", {})
+                    prev_input_length = prev_rec_audit.get("prompt_tokens", 0)
+                    delta = max(0, prev_input - prev_input_length) if prev_input_length > 0 else 0
+                    new_blocks = delta // BLOCK_TOKEN_BUDGET
+                    new_ids = list(
+                        range(global_block_counter, global_block_counter + new_blocks)
+                    )
+                    global_block_counter += new_blocks
+                    cumulative_hash_ids = cumulative_hash_ids + new_ids
+
+                trace_line = {
+                    "chat_id": chat_id,
+                    "parent_chat_id": parent_chat_id,
+                    "timestamp": rec["ts"],
+                    "turn": turn_idx,
+                    "input_length": input_length,
+                    "output_length": output_length,
+                    "type": "chat",
+                    "hash_ids": cumulative_hash_ids,
+                }
+                out.write(json.dumps(trace_line, separators=(",", ":")) + "\n")
+                total_written += 1
+
+    print(f"Converted {total_written} lines from {len(instances)} instances -> {dst}")
+
+
+if __name__ == "__main__":
+    if len(sys.argv) != 3:
+        print(f"Usage: {sys.argv[0]} <input_audit.jsonl> <output_trace.jsonl>")
+        sys.exit(1)
+    convert(Path(sys.argv[1]), Path(sys.argv[2]))

scripts/run_all_experiments.sh

@@ -0,0 +1,73 @@
+#!/bin/bash
+# Run all 3 PD hybrid experiments sequentially
+# Uses 52 sessions / 4,449 requests (10% sample of 497 sessions)
+# Each experiment takes ~30-40 min
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+TRACE="outputs/qwen35-swebench-50sess.jsonl"
+MODEL="/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B"
+OUTPUT="outputs/swebench-exps"
+
+echo "=== Experiment A: pd-disaggregation ==="
+uv run agentic-pd-hybrid benchmark-live \
+  --trace "$TRACE" \
+  --output-root "$OUTPUT" \
+  --mechanism pd-disaggregation \
+  --policy default \
+  --model-path "$MODEL" \
+  --prefill-workers 1 --decode-workers 1 \
+  --prefill-tp-size 4 --decode-tp-size 4 \
+  --prefill-gpu-ids 0,1,2,3 --decode-gpu-ids 4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300
+
+echo "=== Experiment B: pd-colo ==="
+uv run agentic-pd-hybrid benchmark-live \
+  --trace "$TRACE" \
+  --output-root "$OUTPUT" \
+  --mechanism pd-colo \
+  --policy default \
+  --model-path "$MODEL" \
+  --prefill-workers 0 --decode-workers 0 \
+  --direct-workers 2 --direct-tp-size 4 \
+  --direct-gpu-ids 0,1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300
+
+echo "=== Experiment C: kvcache-centric ==="
+uv run agentic-pd-hybrid benchmark-live \
+  --trace "$TRACE" \
+  --output-root "$OUTPUT" \
+  --mechanism kvcache-centric \
+  --policy default \
+  --model-path "$MODEL" \
+  --prefill-workers 1 --decode-workers 1 \
+  --prefill-tp-size 4 --decode-tp-size 4 \
+  --prefill-gpu-ids 0,1,2,3 --decode-gpu-ids 4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 2 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+echo "=== All experiments complete ==="

scripts/run_exp_a_pd_disagg.sh

@@ -0,0 +1,24 @@
+#!/bin/bash
+# Experiment A: pd-disaggregation baseline
+# 1P(GPU 0-3) + 1D(GPU 4-7), TP4, mooncake TCP
+# Full 39K trace from SWE-Bench 500 instances
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/qwen35-swebench-500.jsonl \
+  --output-root outputs/swebench-exps \
+  --mechanism pd-disaggregation \
+  --policy default \
+  --model-path /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B \
+  --prefill-workers 1 --decode-workers 1 \
+  --prefill-tp-size 4 --decode-tp-size 4 \
+  --prefill-gpu-ids 0,1,2,3 --decode-gpu-ids 4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 64 \
+  --timeout-s 900 \
+  --request-timeout-s 300

scripts/run_exp_b1_dp_colo_rr.sh

@@ -0,0 +1,23 @@
+#!/bin/bash
+# Experiment B1: Naive DP colocation — round-robin policy
+# 2 direct workers (GPU 0-3, 4-7), TP4, DP router with round-robin
+# No disaggregation — each worker does prefill+decode locally
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/qwen35-swebench-50sess.jsonl \
+  --output-root outputs/swebench-exps \
+  --mechanism pd-colo \
+  --policy default \
+  --model-path /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B \
+  --prefill-workers 0 --decode-workers 0 \
+  --direct-workers 2 --direct-tp-size 4 \
+  --direct-gpu-ids 0,1,2,3,4,5,6,7 \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300

scripts/run_exp_b2_dp_colo_cache_aware.sh

@@ -0,0 +1,23 @@
+#!/bin/bash
+# Experiment B2: Naive DP colocation — cache-aware (kv-aware) policy
+# 2 direct workers (GPU 0-3, 4-7), TP4, DP router with consistent-hashing
+# Replay kv-aware policy picks the worker with most prefix overlap
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/qwen35-swebench-50sess.jsonl \
+  --output-root outputs/swebench-exps \
+  --mechanism pd-colo \
+  --policy kv-aware \
+  --model-path /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B \
+  --prefill-workers 0 --decode-workers 0 \
+  --direct-workers 2 --direct-tp-size 4 \
+  --direct-gpu-ids 0,1,2,3,4,5,6,7 \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300

scripts/run_exp_b_pd_colo.sh

@@ -0,0 +1,24 @@
+#!/bin/bash
+# Experiment B: pd-colo (direct/colocation)
+# 2 direct workers (GPU 0-3, 4-7), TP4, no router
+# Full 39K trace from SWE-Bench 500 instances
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/qwen35-swebench-500.jsonl \
+  --output-root outputs/swebench-exps \
+  --mechanism pd-colo \
+  --policy default \
+  --model-path /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B \
+  --prefill-workers 0 --decode-workers 0 \
+  --direct-workers 2 --direct-tp-size 4 \
+  --direct-gpu-ids 0,1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 64 \
+  --timeout-s 900 \
+  --request-timeout-s 300

scripts/run_exp_c_kvcache_centric.sh

@@ -0,0 +1,28 @@
+#!/bin/bash
+# Experiment C: kvcache-centric (session-aware PD)
+# 1P(GPU 0-3) + 1D(GPU 4-7), TP4, mooncake TCP
+# Full 39K trace from SWE-Bench 500 instances
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/qwen35-swebench-500.jsonl \
+  --output-root outputs/swebench-exps \
+  --mechanism kvcache-centric \
+  --policy default \
+  --model-path /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B \
+  --prefill-workers 1 --decode-workers 1 \
+  --prefill-tp-size 4 --decode-tp-size 4 \
+  --prefill-gpu-ids 0,1,2,3 --decode-gpu-ids 4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 64 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 2 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction

scripts/smoke_test.sh

@@ -0,0 +1,30 @@
+#!/bin/bash
+# Smoke test: pd-disaggregation with mooncake TCP, 100 requests
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+# Sample a small trace for smoke testing
+uv run agentic-pd-hybrid sample-sessions \
+  --trace outputs/qwen35-swebench-500.jsonl \
+  --output outputs/qwen35-smoke-3sess.jsonl \
+  --session-sample-rate 0.02 \
+  --min-turns 5 \
+  --target-duration-s 300 \
+  --max-requests 100
+
+# Run smoke test
+uv run agentic-pd-hybrid benchmark-live \
+  --trace outputs/qwen35-smoke-3sess.jsonl \
+  --output-root outputs/smoke \
+  --mechanism pd-disaggregation \
+  --policy default \
+  --model-path /mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3.5-35B-A3B \
+  --prefill-workers 1 --decode-workers 1 \
+  --prefill-tp-size 4 --decode-tp-size 4 \
+  --prefill-gpu-ids 0,1,2,3 --decode-gpu-ids 4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300

scripts/sweep_backpressure_smoke.sh

@@ -0,0 +1,114 @@
+#!/usr/bin/env bash
+# Smoke sweep: validate backpressure code change on top of v5 Option D config.
+# Designed to fit in ~3-4h GPU budget (4 runs × ~30-60 min).
+#
+# Usage:
+#   bash scripts/sweep_backpressure_smoke.sh
+#
+# Prerequisites: GPUs available; trace at outputs/qwen35-swebench-50sess.jsonl;
+# model at $MODEL_PATH (default Qwen3-30B-A3B-Instruct-2507).
+set -euo pipefail
+
+REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+cd "$REPO_ROOT"
+
+OUT_ROOT=${OUT_ROOT:-outputs/sweep_backpressure_smoke}
+TRACE=${TRACE:-outputs/qwen35-swebench-50sess.jsonl}
+MODEL=${MODEL:-/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507}
+
+mkdir -p "$OUT_ROOT"
+LOG="$OUT_ROOT/sweep.log"
+echo "[$(date '+%F %T')] Starting backpressure smoke sweep" | tee -a "$LOG"
+echo "  Trace: $TRACE" | tee -a "$LOG"
+echo "  Model: $MODEL" | tee -a "$LOG"
+echo "  Output root: $OUT_ROOT" | tee -a "$LOG"
+
+KVC_COMMON_ARGS=(
+    --trace "$TRACE"
+    --model "$MODEL"
+    --mechanism kvcache-centric
+    --policy kv-aware
+    --kvcache-admission-mode worker
+    --kvcache-seed-min-turn-id 1
+    --kvcache-seed-max-inflight-decode -1
+    --kvcache-prefill-backup-policy release-after-transfer
+    --kvcache-prefill-priority-eviction
+    --prefill-workers 2
+    --decode-workers 6
+    --prefill-gpu-ids 0,1
+    --decode-gpu-ids 2,3,4,5,6,7
+    --transfer-backend mooncake
+    --target-duration-s 2000
+    --session-sample-rate 1.0
+    --min-turns 2
+    --concurrency-limit 32
+)
+
+DP_COMMON_ARGS=(
+    --trace "$TRACE"
+    --model "$MODEL"
+    --mechanism pd-colo
+    --policy kv-aware
+    --direct-workers 8
+    --direct-gpu-ids 0,1,2,3,4,5,6,7
+    --transfer-backend mooncake
+    --target-duration-s 2000
+    --session-sample-rate 1.0
+    --min-turns 2
+    --concurrency-limit 32
+)
+
+run_kvc_baseline_ts10() {
+    local out="$OUT_ROOT/E1_kvc_baseline_ts10"
+    echo "[$(date '+%F %T')] === E1: KVC baseline (no backpressure) time-scale=10 ===" | tee -a "$LOG"
+    python -m agentic_pd_hybrid.cli benchmark-live \
+        "${KVC_COMMON_ARGS[@]}" \
+        --output-root "$out" \
+        --time-scale 10 \
+        2>&1 | tee -a "$LOG"
+}
+
+run_kvc_backpressure_ts10() {
+    local out="$OUT_ROOT/E2_kvc_backpressure_ts10"
+    echo "[$(date '+%F %T')] === E2: KVC + backpressure ON, time-scale=10 ===" | tee -a "$LOG"
+    python -m agentic_pd_hybrid.cli benchmark-live \
+        "${KVC_COMMON_ARGS[@]}" \
+        --output-root "$out" \
+        --time-scale 10 \
+        --enable-backpressure \
+        --backpressure-max-pause-s 2.0 \
+        2>&1 | tee -a "$LOG"
+}
+
+run_kvc_backpressure_ts1() {
+    local out="$OUT_ROOT/E3_kvc_backpressure_ts1_short"
+    echo "[$(date '+%F %T')] === E3: KVC + backpressure ON, time-scale=1, FIRST 1000 reqs ===" | tee -a "$LOG"
+    python -m agentic_pd_hybrid.cli benchmark-live \
+        "${KVC_COMMON_ARGS[@]}" \
+        --output-root "$out" \
+        --time-scale 1 \
+        --enable-backpressure \
+        --backpressure-max-pause-s 2.0 \
+        --target-duration-s 1800 \
+        2>&1 | tee -a "$LOG"
+}
+
+run_dp_baseline_ts1() {
+    local out="$OUT_ROOT/E4_dp_ts1_short"
+    echo "[$(date '+%F %T')] === E4: 8-way DP cache-aware, time-scale=1, FIRST 1000 reqs ===" | tee -a "$LOG"
+    python -m agentic_pd_hybrid.cli benchmark-live \
+        "${DP_COMMON_ARGS[@]}" \
+        --output-root "$out" \
+        --time-scale 1 \
+        --target-duration-s 1800 \
+        2>&1 | tee -a "$LOG"
+}
+
+# Sequence — add/remove as fits the budget.
+run_kvc_baseline_ts10
+run_kvc_backpressure_ts10
+run_kvc_backpressure_ts1
+run_dp_baseline_ts1
+
+echo "[$(date '+%F %T')] === sweep DONE ===" | tee -a "$LOG"
+echo "Run analysis with: python scripts/analysis/analyze_backpressure_smoke.py $OUT_ROOT" | tee -a "$LOG"

scripts/sweep_kvc_qwen3_30b.sh

@@ -0,0 +1,60 @@
+#!/bin/bash
+# KVC admission control parameter sweep on Qwen3-30B
+# 5 experiments, ~35 min each, ~3 hours total
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-exps
+VENV_PYTHON=.venv/bin/python
+
+run_kvc() {
+  local label=$1
+  local inflight=$2
+  local min_turn=$3
+
+  echo "=== [$label] inflight=$inflight min_turn=$min_turn === $(date)"
+  PYTHONPATH=src:third_party/sglang/python \
+  $VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+    --trace $TRACE \
+    --output-root $OUTPUT \
+    --mechanism kvcache-centric \
+    --policy default \
+    --model-path $MODEL \
+    --prefill-workers 1 --decode-workers 1 \
+    --prefill-tp-size 4 --decode-tp-size 4 \
+    --prefill-gpu-ids 0,1,2,3 --decode-gpu-ids 4,5,6,7 \
+    --transfer-backend mooncake \
+    --gpu-budget 8 \
+    --time-scale 10 \
+    --session-sample-rate 1.0 \
+    --target-duration-s 100000 \
+    --concurrency-limit 32 \
+    --timeout-s 900 \
+    --request-timeout-s 300 \
+    --kvcache-admission-mode worker \
+    --kvcache-seed-min-turn-id $min_turn \
+    --kvcache-seed-max-inflight-decode $inflight \
+    --kvcache-prefill-backup-policy release-after-transfer \
+    --kvcache-prefill-priority-eviction
+  echo "=== [$label] DONE === $(date)"
+  echo ""
+}
+
+# C1: inflight=8, min-turn=2
+run_kvc "C1" 8 2
+
+# C2: inflight=16, min-turn=2
+run_kvc "C2" 16 2
+
+# C3: inflight=-1 (disabled), min-turn=2
+run_kvc "C3" -1 2
+
+# C4: inflight=8, min-turn=1
+run_kvc "C4" 8 1
+
+# C5: inflight=-1 (disabled), min-turn=1
+run_kvc "C5" -1 1
+
+echo "=== ALL SWEEP EXPERIMENTS DONE === $(date)"

scripts/sweep_tp1_configs.sh

@@ -0,0 +1,133 @@
+#!/bin/bash
+# TP1 configuration sweep: 8-way DP, 1P7D KVC, 2P6D KVC
+# Qwen3-30B-A3B TP=1, single GPU per worker
+# Most aggressive KVC admission: inflight=-1 (off), seed-min-turn=1
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-exps
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    # Also copy summary to a named file for easy access
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    log "Saved to $OUTPUT/${label}_summary.json"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting TP1 configuration sweep"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+
+########################################
+# Experiment 1: 8-way DP cache-aware
+########################################
+log ""
+log "=== [EXP1] 8-way DP cache-aware (8 direct × TP1) ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism pd-colo \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 0 --decode-workers 0 \
+  --direct-workers 8 --direct-tp-size 1 \
+  --direct-gpu-ids 0,1,2,3,4,5,6,7 \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300
+
+# Find latest run dir for this experiment
+EXP1_DIR=$(ls -td $OUTPUT/pd-colo-kv-aware-*/ 2>/dev/null | head -1)
+save_result "exp1_8way_dp_cache_aware" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 1P + 7D KVC (most aggressive)
+########################################
+log ""
+log "=== [EXP2] 1P7D KVC (inflight=off, min-turn=1) ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy default \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_1p7d_kvc_aggressive" "$EXP2_DIR"
+
+########################################
+# Experiment 3: 2P + 6D KVC (most aggressive)
+########################################
+log ""
+log "=== [EXP3] 2P6D KVC (inflight=off, min-turn=1) ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy default \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP3_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp3_2p6d_kvc_aggressive" "$EXP3_DIR"
+
+########################################
+log ""
+log "=== ALL TP1 SWEEP EXPERIMENTS DONE ==="

scripts/sweep_tp1_v2_fixed.sh

@@ -0,0 +1,131 @@
+#!/bin/bash
+# TP1 configuration sweep v2 — after session_params fix + audit fields
+# Qwen3-30B-A3B TP=1, single GPU per worker
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v2-fixed
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    log "Saved to $OUTPUT/${label}_summary.json + ${label}_metrics.jsonl"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting TP1 v2 sweep (session_params fix + audit fields)"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+
+########################################
+# Experiment 1: 8-way DP cache-aware
+########################################
+log ""
+log "=== [EXP1] 8-way DP cache-aware (8 direct × TP1) ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism pd-colo \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 0 --decode-workers 0 \
+  --direct-workers 8 --direct-tp-size 1 \
+  --direct-gpu-ids 0,1,2,3,4,5,6,7 \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300
+
+EXP1_DIR=$(ls -td $OUTPUT/pd-colo-kv-aware-*/ 2>/dev/null | head -1)
+save_result "exp1_8way_dp_cache_aware" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 1P + 7D KVC (aggressive)
+########################################
+log ""
+log "=== [EXP2] 1P7D KVC (inflight=off, min-turn=1) ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy default \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_1p7d_kvc_aggressive" "$EXP2_DIR"
+
+########################################
+# Experiment 3: 2P + 6D KVC (aggressive)
+########################################
+log ""
+log "=== [EXP3] 2P6D KVC (inflight=off, min-turn=1) ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy default \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP3_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp3_2p6d_kvc_aggressive" "$EXP3_DIR"
+
+########################################
+log ""
+log "=== ALL TP1 V2 SWEEP EXPERIMENTS DONE ==="

scripts/sweep_tp1_v3_kvaware.sh

@@ -0,0 +1,108 @@
+#!/bin/bash
+# TP1 v3 sweep — KVC with kv-aware policy (fix routing mismatch)
+# v2 used --policy default for KVC experiments, causing session routing
+# mismatch: replay round-robin ≠ router round-robin → "session not found".
+# v3 uses --policy kv-aware for KVC to ensure session affinity.
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v3-kvaware
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    log "Saved to $OUTPUT/${label}_summary.json + ${label}_metrics.jsonl"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting TP1 v3 sweep (KVC with kv-aware policy)"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Key change: --policy kv-aware for KVC (was --policy default in v2)"
+
+########################################
+# Experiment 1: 1P + 7D KVC kv-aware
+########################################
+log ""
+log "=== [EXP1] 1P7D KVC kv-aware ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP1_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp1_1p7d_kvc_kvaware" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 2P + 6D KVC kv-aware
+########################################
+log ""
+log "=== [EXP2] 2P6D KVC kv-aware ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_2p6d_kvc_kvaware" "$EXP2_DIR"
+
+########################################
+log ""
+log "=== ALL TP1 V3 SWEEP EXPERIMENTS DONE ==="

scripts/sweep_tp1_v4_cap16.sh

@@ -0,0 +1,108 @@
+#!/bin/bash
+# TP1 v4 sweep — KVC with kv-aware policy + soft_cap raised from 4 to 16
+# v3 (kv-aware) fixed routing but session-cap fallback still dominated 52-65%
+# of requests. Hardcoded min(4, ...) in _decode_session_soft_cap was the
+# bottleneck — only 4*7=28 session slots for 52 trace sessions.
+# v4 raises the cap to 16 (4*7=28 -> 16*7=112 slots).
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v4-cap16
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    log "Saved to $OUTPUT/${label}_summary.json + ${label}_metrics.jsonl"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting TP1 v4 sweep (KVC kv-aware, session soft_cap raised 4->16)"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Key change: _decode_session_soft_cap now min(16, ...) instead of min(4, ...)"
+
+########################################
+# Experiment 1: 1P + 7D KVC kv-aware (cap=16)
+########################################
+log ""
+log "=== [EXP1] 1P7D KVC kv-aware cap=16 ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP1_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp1_1p7d_kvc_cap16" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 2P + 6D KVC kv-aware (cap=16)
+########################################
+log ""
+log "=== [EXP2] 2P6D KVC kv-aware cap=16 ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_2p6d_kvc_cap16" "$EXP2_DIR"
+
+log ""
+log "=== ALL TP1 V4 SWEEP EXPERIMENTS DONE ==="

scripts/sweep_tp1_v5_baseline_rerun_exp2.sh

@@ -0,0 +1,89 @@
+#!/bin/bash
+# P0: Re-run v5 baseline EXP2 (2P6D) three times to establish whether
+# errors=9 is a stable property of the v5 config or single-run variance.
+# Critic of V5_PROFILE_INVESTIGATION_ZH.md flagged that the 415 errors in
+# v5+profile EXP2 may have been polling-induced. We need 3 baseline runs
+# (no polling, identical config to original v5) to test reproducibility.
+#
+# Output:
+#   outputs/qwen3-30b-tp1-v5-optD-baseline-rerun/
+#     ├── exp2_2p6d_run{1,2,3}_summary.json
+#     ├── exp2_2p6d_run{1,2,3}_metrics.jsonl
+#     └── kvcache-centric-...<ts>/   (one per run)
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v5-optD-baseline-rerun
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+run_exp2() {
+  local run_idx=$1
+  local label="exp2_2p6d_run${run_idx}"
+  log ""
+  log "=== [RUN ${run_idx}/3] EXP2 2P6D KVC kv-aware Option D (no polling) ==="
+  PYTHONPATH=src:third_party/sglang/python \
+  $VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+    --trace $TRACE \
+    --output-root $OUTPUT \
+    --mechanism kvcache-centric \
+    --policy kv-aware \
+    --model-path $MODEL \
+    --prefill-workers 2 --decode-workers 6 \
+    --prefill-tp-size 1 --decode-tp-size 1 \
+    --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+    --transfer-backend mooncake \
+    --gpu-budget 8 \
+    --time-scale 10 \
+    --session-sample-rate 1.0 \
+    --target-duration-s 100000 \
+    --concurrency-limit 32 \
+    --timeout-s 900 \
+    --request-timeout-s 300 \
+    --kvcache-admission-mode worker \
+    --kvcache-seed-min-turn-id 1 \
+    --kvcache-seed-max-inflight-decode -1 \
+    --kvcache-prefill-backup-policy release-after-transfer \
+    --kvcache-prefill-priority-eviction
+
+  local run_dir=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+  log "=== [RUN ${run_idx}/3] $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    local errs=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('error_count',0))")
+    log "  errors = $errs (baseline reference = 9)"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+  else
+    log "WARNING: no summary file in $run_dir"
+  fi
+}
+
+log "=== P0: v5 baseline EXP2 reproducibility test (3 runs, no polling) ==="
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Goal: confirm whether errors=9 in v5 baseline EXP2 is reproducible"
+log "      (v5+profile saw 415 errors; we need to know if polling was causal)"
+
+for i in 1 2 3; do
+  run_exp2 $i
+done
+
+log ""
+log "=== P0 SUMMARY: errors per run ==="
+for i in 1 2 3; do
+  if [ -f "$OUTPUT/exp2_2p6d_run${i}_summary.json" ]; then
+    e=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/exp2_2p6d_run${i}_summary.json')); print(d.get('error_count',0))")
+    log "  run ${i}: errors = $e"
+  fi
+done
+log "=== P0 ALL DONE ==="

scripts/sweep_tp1_v5_optD.sh

@@ -0,0 +1,114 @@
+#!/bin/bash
+# TP1 v5 sweep — Option D: D-side admission for seed/reseed.
+#
+# v4 (cap=16) still saw 35% session-cap fallback because the local soft_cap
+# evaluates min(16, usable_capacity_tokens / target_tokens) and target_tokens
+# (= input + output) is 50-100K in agentic workloads, giving cap = 1-2.
+#
+# v5 makes worker admission_mode authoritative for ALL admission decisions
+# (direct_append AND seed/reseed). Replay calls D's
+# /session_cache/admit_direct_append with mode={direct_append|seed} and
+# defers to D's KV pool availability + LRU eviction. Replay's local
+# _decode_session_soft_cap is bypassed entirely under worker mode.
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v5-optD
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    log "Saved to $OUTPUT/${label}_summary.json + ${label}_metrics.jsonl"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting TP1 v5 sweep (Option D: D-side seed admission)"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Key change: worker admission_mode now drives seed/reseed via D's admit endpoint"
+
+########################################
+# Experiment 1: 1P + 7D KVC kv-aware Option D
+########################################
+log ""
+log "=== [EXP1] 1P7D KVC kv-aware Option D ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP1_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp1_1p7d_kvc_optD" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 2P + 6D KVC kv-aware Option D
+########################################
+log ""
+log "=== [EXP2] 2P6D KVC kv-aware Option D ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_2p6d_kvc_optD" "$EXP2_DIR"
+
+log ""
+log "=== ALL TP1 V5 SWEEP EXPERIMENTS DONE ==="

scripts/sweep_tp1_v5_optD_profile.sh

@@ -0,0 +1,125 @@
+#!/bin/bash
+# TP1 v5 + profiling — re-run the v5 (Option D) config with the new
+# d-pool-timeseries poller enabled, so we can attribute each session-cap
+# fallback to actual D KV pool occupancy (held vs available vs idle-evictable
+# vs prefill-backup) instead of guessing.
+#
+# Output:
+#   outputs/qwen3-30b-tp1-v5-optD-profile/
+#     ├── kvcache-centric-kv-aware-worker-admission-<ts>/
+#     │   ├── request-metrics.jsonl
+#     │   ├── request-metrics.jsonl.summary.json
+#     │   └── d-pool-timeseries.jsonl   ← NEW (1Hz P/D /server_info snapshots)
+#     ├── exp1_1p7d_kvc_optD_profile_metrics.jsonl
+#     └── exp2_2p6d_kvc_optD_profile_metrics.jsonl
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v5-optD-profile
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+POLL_INTERVAL=1.0
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    if [ -f "$run_dir/d-pool-timeseries.jsonl" ]; then
+      cp "$run_dir/d-pool-timeseries.jsonl" "$OUTPUT/${label}_pool_timeseries.jsonl"
+      log "Pool timeseries: $(wc -l < $OUTPUT/${label}_pool_timeseries.jsonl) rows"
+    else
+      log "WARNING: no d-pool-timeseries.jsonl produced"
+    fi
+    log "Saved to $OUTPUT/${label}_summary.json + ${label}_metrics.jsonl + ${label}_pool_timeseries.jsonl"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting TP1 v5 + profile sweep (Option D + ${POLL_INTERVAL}s pool polling)"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Profiling: --pool-poll-interval-s $POLL_INTERVAL (writes d-pool-timeseries.jsonl)"
+
+########################################
+# Experiment 1: 1P + 7D KVC kv-aware Option D + profile
+########################################
+log ""
+log "=== [EXP1] 1P7D KVC kv-aware Option D + profile ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction \
+  --pool-poll-interval-s $POLL_INTERVAL
+
+EXP1_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp1_1p7d_kvc_optD_profile" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 2P + 6D KVC kv-aware Option D + profile
+########################################
+log ""
+log "=== [EXP2] 2P6D KVC kv-aware Option D + profile ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction \
+  --pool-poll-interval-s $POLL_INTERVAL
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_2p6d_kvc_optD_profile" "$EXP2_DIR"
+
+log ""
+log "=== ALL TP1 V5+PROFILE EXPERIMENTS DONE ==="

scripts/sweep_tp1_v6_p1_profile.sh

@@ -0,0 +1,129 @@
+#!/bin/bash
+# v6 P1: re-run the v5 (Option D) config with the pool_breakdown instrument
+# (commit 4978c0d) so d-pool-timeseries.jsonl carries radix_protected /
+# slot_private / running_batch / {transfer,prealloc,retracted}_queue tokens.
+#
+# This is the same config as scripts/sweep_tp1_v5_optD_profile.sh but writes
+# to a separate output dir, leaving the pre-instrument v5+profile run intact
+# for before/after comparison.
+#
+# Output:
+#   outputs/qwen3-30b-tp1-v6-p1-profile/
+#     ├── kvcache-centric-kv-aware-worker-admission-<ts>/
+#     │   ├── request-metrics.jsonl
+#     │   ├── request-metrics.jsonl.summary.json
+#     │   └── d-pool-timeseries.jsonl   ← now with pool_breakdown fields
+#     ├── exp{1,2}_*_metrics.jsonl
+#     └── exp{1,2}_*_pool_timeseries.jsonl
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-v6-p1-profile
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+POLL_INTERVAL=1.0
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+save_result() {
+  local label=$1
+  local run_dir=$2
+  log "=== $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    log "Summary:"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    if [ -f "$run_dir/d-pool-timeseries.jsonl" ]; then
+      cp "$run_dir/d-pool-timeseries.jsonl" "$OUTPUT/${label}_pool_timeseries.jsonl"
+      log "Pool timeseries: $(wc -l < $OUTPUT/${label}_pool_timeseries.jsonl) rows"
+    else
+      log "WARNING: no d-pool-timeseries.jsonl produced"
+    fi
+    log "Saved to $OUTPUT/${label}_summary.json + ${label}_metrics.jsonl + ${label}_pool_timeseries.jsonl"
+  else
+    log "WARNING: No summary file found in $run_dir"
+  fi
+}
+
+log "Starting v6 P1 sweep (v5 Option D config + ${POLL_INTERVAL}s pool polling + pool_breakdown)"
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Goal: capture pool_breakdown fields (radix_protected / slot_private / running_batch / queues)"
+log "      to decompose 'other' on the v5 baseline workload"
+
+########################################
+# Experiment 1: 1P + 7D KVC kv-aware Option D + profile
+########################################
+log ""
+log "=== [EXP1] 1P7D KVC kv-aware Option D + profile ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 7 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction \
+  --pool-poll-interval-s $POLL_INTERVAL
+
+EXP1_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp1_1p7d_kvc_v6_p1" "$EXP1_DIR"
+
+########################################
+# Experiment 2: 2P + 6D KVC kv-aware Option D + profile
+########################################
+log ""
+log "=== [EXP2] 2P6D KVC kv-aware Option D + profile ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 2 --decode-workers 6 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0,1 --decode-gpu-ids 2,3,4,5,6,7 \
+  --transfer-backend mooncake \
+  --gpu-budget 8 \
+  --time-scale 10 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction \
+  --pool-poll-interval-s $POLL_INTERVAL
+
+EXP2_DIR=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+save_result "exp2_2p6d_kvc_v6_p1" "$EXP2_DIR"
+
+log ""
+log "=== ALL v6 P1 EXPERIMENTS DONE ==="

scripts/sweep_ts1_kvc_n3_plus_dp.sh

@@ -0,0 +1,146 @@
+#!/bin/bash
+# Time-scale=1 validation sweep, downscaled to 4 GPUs:
+#   - KVC v5 1P3D × N=3   (new data, validates §1/§2 structural claims at real timing)
+#   - 4-way DP cache-aware × 1 (sanity baseline at same scale + ts=1)
+#
+# Goal: per docs/AGENTIC_FIT_ANALYSIS_ZH.md §7 / TEAM_REPORT §2.6 — all v3-v6 KVC
+# data was at time-scale=10 (inter-turn gap p50 = 0.25s, vs real 2.5s). This run
+# tests whether the gap structurally reverses any conclusion.
+#
+# CONFIG NOTE: Original experiments used 8 GPUs (2P6D / 8-way DP). This host has
+# only 4 H100s available, so we downscale proportionally to 1P3D / 4-way DP.
+# Cross-compare against existing 2P6D ts=10 data is confounded by *both*
+# time-scale and capacity. Internal comparison (1P3D KVC vs 4DP) at ts=1 is the
+# clean signal. §5 (P-side imbalance) is NOT testable here — only 1 P.
+#
+# Capacity ratio: 3D × ~92K tok = 276K KV pool vs 52 sessions × ~50K peak input
+#   working set ≈ 1.5M → ~5.4× overload (vs 2.7× in original 2P6D).
+#   Pressure is HIGHER than original; partly offset by ts=1 letting D drain between turns.
+#
+# Output:
+#   outputs/qwen3-30b-tp1-ts1-validation/
+#     ├── kvc_1p3d_run{1,2,3}_summary.json
+#     ├── kvc_1p3d_run{1,2,3}_metrics.jsonl
+#     ├── dp4_summary.json
+#     ├── dp4_metrics.jsonl
+#     └── kvcache-centric-... / pd-colo-kv-aware-...    (raw run dirs)
+#
+# Estimated GPU time: KVC ts=1 ≈ 100-180 min/run × 3 = 5-9h
+#                     DP ts=1  ≈ 100-120 min × 1     = ~2h
+#                     Total                           = 7-11h
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-ts1-validation
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() {
+  echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE
+}
+
+run_kvc_1p3d() {
+  local run_idx=$1
+  local label="kvc_1p3d_run${run_idx}"
+  log ""
+  log "=== [KVC ${run_idx}/3] 1P3D KVC kv-aware Option D, time-scale=1 ==="
+  PYTHONPATH=src:third_party/sglang/python \
+  $VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+    --trace $TRACE \
+    --output-root $OUTPUT \
+    --mechanism kvcache-centric \
+    --policy kv-aware \
+    --model-path $MODEL \
+    --prefill-workers 1 --decode-workers 3 \
+    --prefill-tp-size 1 --decode-tp-size 1 \
+    --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3 \
+    --transfer-backend mooncake \
+    --gpu-budget 4 \
+    --time-scale 1 \
+    --session-sample-rate 1.0 \
+    --target-duration-s 100000 \
+    --concurrency-limit 32 \
+    --timeout-s 900 \
+    --request-timeout-s 300 \
+    --kvcache-admission-mode worker \
+    --kvcache-seed-min-turn-id 1 \
+    --kvcache-seed-max-inflight-decode -1 \
+    --kvcache-prefill-backup-policy release-after-transfer \
+    --kvcache-prefill-priority-eviction
+
+  local run_dir=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+  log "=== [KVC ${run_idx}/3] $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    local errs=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('error_count',0))")
+    log "  errors = $errs"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+  else
+    log "WARNING: no summary file in $run_dir"
+  fi
+}
+
+run_dp4_sanity() {
+  local label="dp4"
+  log ""
+  log "=== [DP] 4-way DP cache-aware sanity, time-scale=1 ==="
+  PYTHONPATH=src:third_party/sglang/python \
+  $VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+    --trace $TRACE \
+    --output-root $OUTPUT \
+    --mechanism pd-colo \
+    --policy kv-aware \
+    --model-path $MODEL \
+    --prefill-workers 0 --decode-workers 0 \
+    --direct-workers 4 --direct-tp-size 1 \
+    --direct-gpu-ids 0,1,2,3 \
+    --gpu-budget 4 \
+    --time-scale 1 \
+    --session-sample-rate 1.0 \
+    --target-duration-s 100000 \
+    --concurrency-limit 32 \
+    --timeout-s 900 \
+    --request-timeout-s 300
+
+  local run_dir=$(ls -td $OUTPUT/pd-colo-kv-aware-*/ 2>/dev/null | head -1)
+  log "=== [DP] $label COMPLETED ==="
+  if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+    cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+    cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+    local errs=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('error_count',0))")
+    log "  errors = $errs"
+    cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+    echo "" >> $RESULTS_FILE
+  else
+    log "WARNING: no summary file in $run_dir"
+  fi
+}
+
+log "=== TS=1 VALIDATION (4-GPU): KVC 1P3D × N=3 + 4DP × 1 ==="
+log "Model: $MODEL"
+log "Trace: $TRACE (4449 requests, 52 sessions)"
+log "Goal: validate whether ts=10 was the main distortion in v3-v6 KVC vs DP"
+
+# KVC × 3 first (the new data we need); DP last (cheaper sanity at end)
+for i in 1 2 3; do
+  run_kvc_1p3d $i
+done
+
+run_dp4_sanity
+
+log ""
+log "=== TS=1 SUMMARY ==="
+for label in kvc_1p3d_run1 kvc_1p3d_run2 kvc_1p3d_run3 dp4; do
+  if [ -f "$OUTPUT/${label}_summary.json" ]; then
+    e=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('error_count',0))")
+    p50=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('latency_stats_s',{}).get('p50','n/a'))")
+    log "  ${label}: errors=$e  lat_p50=${p50}s"
+  fi
+done
+log "=== TS=1 ALL DONE ==="

scripts/sweep_ts1_migration_v1.sh

@@ -0,0 +1,65 @@
+#!/bin/bash
+# Migration v1 validation: KVC 1P3D ts=1 with --kvcache-migration-reject-threshold=3
+# Compare against baseline outputs/qwen3-30b-tp1-ts1-validation/kvc_1p3d_run{1,2,3}
+# (all of which had no migration — runs were structurally identical).
+#
+# Goal: verify §1 fix changes the categorical outcome — direct-to-D % up,
+# fallback-session-not-resident % down, lat mean down.
+#
+# ts=1 is deterministic at the categorical level, so N=1 is sufficient
+# (TEAM_REPORT §2.8 revised).
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-ts1-migration-v1
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() { echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE; }
+
+log "=== TS=1 MIGRATION v1: KVC 1P3D --kvcache-migration-reject-threshold=3 ==="
+log "Baseline reference: outputs/qwen3-30b-tp1-ts1-validation/kvc_1p3d_run1 (errors=5, lat mean=1.574s, direct-to-D=42.8%)"
+
+label=kvc_1p3d_migration_run1
+log ""
+log "=== [migration v1] starting ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 3 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3 \
+  --transfer-backend mooncake \
+  --gpu-budget 4 \
+  --time-scale 1 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction \
+  --kvcache-migration-reject-threshold 3
+
+run_dir=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+log "=== [migration v1] $label COMPLETED ==="
+if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+  cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+  cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+  errs=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('error_count',0))")
+  p50=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('latency_stats_s',{}).get('p50',0))")
+  log "  errors=$errs lat_p50=${p50}s"
+  cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+fi
+log "=== migration v1 DONE ==="

scripts/sweep_ts1_migration_v2.sh

@@ -0,0 +1,76 @@
+#!/bin/bash
+# Migration v2 validation: KVC 1P3D ts=1 with BOTH:
+#   (1) reset-on-success blacklist decay (replay.py code change)
+#   (2) --kvcache-direct-max-uncached-tokens 8192 (was 2048 default)
+#
+# v1 results (kvc_1p3d_migration_run1) showed:
+#   - lat mean WORSE +11.7%, TTFT mean WORSE +71.3% — thrashing tax
+#   - direct-to-D rate UP +10.5pp (42.8 → 53.3%)
+#   - Fallback breakdown surprise: 41.3% are 'real-large-append' (>2048 tok),
+#     NOT 'session-not-resident' as we hypothesized
+#
+# v2 design (REFACTOR_PLAN_V1 + MIGRATION_V1_FINDINGS):
+#   (1) reset-on-success: clear (sess,D) reject counter on successful direct-to-D
+#       — eliminates blacklist-permanence bug → kills thrashing
+#   (2) bump direct-append threshold 2048 → 8192: lets more large-append turns
+#       go direct-to-D instead of fall through to seed (which often rejects)
+set -euo pipefail
+cd "$(dirname "$0")/.."
+
+MODEL=/mnt/kzlin/workflow/pd-hybrid/simm-swe-bench/models/Qwen3-30B-A3B-Instruct-2507
+TRACE=outputs/qwen35-swebench-50sess.jsonl
+OUTPUT=outputs/qwen3-30b-tp1-ts1-migration-v2
+VENV_PYTHON=.venv/bin/python
+RESULTS_FILE=$OUTPUT/sweep_results.txt
+
+mkdir -p $OUTPUT
+
+log() { echo "[$(date '+%Y-%m-%d %H:%M:%S')] $*" | tee -a $RESULTS_FILE; }
+
+log "=== TS=1 MIGRATION v2: reset-on-success + threshold=8192 ==="
+log "Baselines:"
+log "  baseline (no migration):        kvc_1p3d_run1 errors=5 lat_p50=0.811s ttft_p50=0.124s direct=42.8%"
+log "  v1 (migration permanent):       kvc_1p3d_migration_run1 errors=6 lat_p50=0.773s ttft_p50=0.057s direct=53.3% lat_mean=1.758s"
+log "  4DP ts=1:                       errors=0 lat_p50=0.659s ttft_p50=0.090s lat_mean=1.443s"
+log "Goal: kill thrashing tax (lat_mean ≤ 1.5s, p99 ≤ 9s) while preserving v1's direct-to-D gains."
+
+label=kvc_1p3d_migration_v2_run1
+log ""
+log "=== [migration v2] starting ==="
+PYTHONPATH=src:third_party/sglang/python \
+$VENV_PYTHON -m agentic_pd_hybrid.cli benchmark-live \
+  --trace $TRACE \
+  --output-root $OUTPUT \
+  --mechanism kvcache-centric \
+  --policy kv-aware \
+  --model-path $MODEL \
+  --prefill-workers 1 --decode-workers 3 \
+  --prefill-tp-size 1 --decode-tp-size 1 \
+  --prefill-gpu-ids 0 --decode-gpu-ids 1,2,3 \
+  --transfer-backend mooncake \
+  --gpu-budget 4 \
+  --time-scale 1 \
+  --session-sample-rate 1.0 \
+  --target-duration-s 100000 \
+  --concurrency-limit 32 \
+  --timeout-s 900 \
+  --request-timeout-s 300 \
+  --kvcache-admission-mode worker \
+  --kvcache-seed-min-turn-id 1 \
+  --kvcache-seed-max-inflight-decode -1 \
+  --kvcache-prefill-backup-policy release-after-transfer \
+  --kvcache-prefill-priority-eviction \
+  --kvcache-migration-reject-threshold 3 \
+  --kvcache-direct-max-uncached-tokens 8192
+
+run_dir=$(ls -td $OUTPUT/kvcache-centric-*/ 2>/dev/null | head -1)
+log "=== [migration v2] $label COMPLETED ==="
+if [ -f "$run_dir/request-metrics.jsonl.summary.json" ]; then
+  cp "$run_dir/request-metrics.jsonl.summary.json" "$OUTPUT/${label}_summary.json"
+  cp "$run_dir/request-metrics.jsonl" "$OUTPUT/${label}_metrics.jsonl"
+  errs=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('error_count',0))")
+  p50=$($VENV_PYTHON -c "import json; d=json.load(open('$OUTPUT/${label}_summary.json')); print(d.get('latency_stats_s',{}).get('p50',0))")
+  log "  errors=$errs lat_p50=${p50}s"
+  cat "$run_dir/request-metrics.jsonl.summary.json" >> $RESULTS_FILE
+fi
+log "=== migration v2 DONE ==="

src/agentic_pd_hybrid/__init__.py

@@ -0,0 +1,12 @@
+"""Agentic PD hybrid prototype."""
+
+__all__ = [
+    "cli",
+    "launcher",
+    "metrics",
+    "microbench",
+    "policies",
+    "replay",
+    "topology",
+    "trace",
+]

src/agentic_pd_hybrid/benchmark.py

@@ -0,0 +1,315 @@
+from __future__ import annotations
+
+import asyncio
+import json
+import signal
+from dataclasses import asdict, dataclass, replace
+from datetime import UTC, datetime
+from pathlib import Path
+
+from agentic_pd_hybrid.replay import ReplayConfig, replay_trace
+from agentic_pd_hybrid.sampling import SessionSampleConfig, sample_trace_sessions
+from agentic_pd_hybrid.stack import ManagedPdStack, launch_pd_stack
+from agentic_pd_hybrid.topology import SingleNodeTopology
+
+
+@dataclass(frozen=True)
+class BenchmarkConfig:
+    trace_path: Path
+    output_root: Path
+    topology: SingleNodeTopology
+    policy_name: str
+    mechanism_name: str = "pd-disaggregation"
+    target_duration_s: float = 600.0
+    start_time_s: float = 0.0
+    session_sample_rate: float = 0.01
+    min_turns: int = 1
+    time_scale: float = 1.0
+    concurrency_limit: int = 32
+    timeout_s: float = 1200.0
+    request_timeout_s: float | None = None
+    stream: bool = True
+    stream_idle_timeout_s: float | None = 900.0
+    kvcache_direct_max_uncached_tokens: int = 2048
+    kvcache_admission_mode: str = "router"
+    kvcache_seed_max_resident_tokens: int | None = None
+    kvcache_seed_max_output_tokens: int | None = None
+    kvcache_seed_min_turn_id: int = 1
+    kvcache_seed_only_multiturn_sessions: bool = False
+    kvcache_prefill_backup_policy: str = "release-after-transfer"
+    kvcache_seed_max_inflight_decode: int | None = 3
+    kvcache_seed_max_decode_transfer_queue_reqs: int | None = None
+    kvcache_direct_max_decode_transfer_queue_reqs: int | None = None
+    kvcache_prefill_priority_eviction: bool = False
+    kvcache_prefill_direct_priority: int = -100
+    kvcache_prefill_normal_priority: int = 100
+    pool_poll_interval_s: float = 0.0
+    pool_poll_include_sessions: bool = True
+    enable_backpressure: bool = False
+    backpressure_max_pause_s: float = 2.0
+    kvcache_migration_reject_threshold: int = 3
+    sample_profile: str = "default"
+    min_initial_input_tokens: int | None = None
+    max_initial_input_tokens: int | None = None
+    max_append_input_tokens: int | None = None
+    max_output_tokens: int | None = None
+    min_overlap_ratio: float | None = None
+    launch_stack: bool = True
+
+
+@dataclass(frozen=True)
+class BenchmarkArtifacts:
+    run_dir: Path
+    sampled_trace_path: Path
+    metrics_path: Path
+    summary_path: Path
+    benchmark_config_path: Path
+
+
+def run_live_benchmark(config: BenchmarkConfig) -> BenchmarkArtifacts:
+    run_id = datetime.now(UTC).strftime("%Y%m%dT%H%M%SZ")
+    run_label = f"{config.mechanism_name}-{config.policy_name}"
+    if config.mechanism_name == "kvcache-centric":
+        run_label = f"{run_label}-{config.kvcache_admission_mode}-admission"
+    run_dir = config.output_root / f"{run_label}-{run_id}"
+    run_dir.mkdir(parents=True, exist_ok=True)
+
+    topology = config.topology
+    if config.mechanism_name == "kvcache-centric":
+        prefill_extra_server_args = topology.prefill_extra_server_args + (
+            "--enable-streaming-session",
+        )
+        if config.kvcache_prefill_priority_eviction:
+            prefill_extra_server_args = prefill_extra_server_args + (
+                "--radix-eviction-policy",
+                "priority",
+            )
+        topology = replace(
+            topology,
+            prefill_extra_server_args=prefill_extra_server_args,
+            decode_extra_server_args=topology.decode_extra_server_args
+            + (
+                "--enable-streaming-session",
+                "--disaggregation-decode-allow-local-prefill",
+            ),
+        )
+
+    sampled_trace_path = run_dir / "sampled-trace.jsonl"
+    sample_summary = sample_trace_sessions(
+        SessionSampleConfig(
+            trace_path=config.trace_path,
+            output_path=sampled_trace_path,
+            target_duration_s=config.target_duration_s,
+            start_time_s=config.start_time_s,
+            session_sample_rate=config.session_sample_rate,
+            min_turns=config.min_turns,
+            profile=config.sample_profile,  # type: ignore[arg-type]
+            min_initial_input_tokens=config.min_initial_input_tokens,
+            max_initial_input_tokens=config.max_initial_input_tokens,
+            max_append_input_tokens=config.max_append_input_tokens,
+            max_output_tokens=config.max_output_tokens,
+            min_overlap_ratio=config.min_overlap_ratio,
+        )
+    )
+
+    stack: ManagedPdStack | None = None
+    previous_sigint = signal.getsignal(signal.SIGINT)
+    previous_sigterm = signal.getsignal(signal.SIGTERM)
+
+    def _handle_termination(signum, _frame) -> None:
+        if stack is not None:
+            stack.stop()
+        raise SystemExit(128 + signum)
+
+    try:
+        signal.signal(signal.SIGINT, _handle_termination)
+        signal.signal(signal.SIGTERM, _handle_termination)
+        _mechanisms_with_router = {"pd-disaggregation", "kvcache-centric", "pd-colo"}
+        _naive_dp = config.mechanism_name == "pd-colo"
+        if config.launch_stack:
+            stack = launch_pd_stack(
+                topology=topology,
+                run_dir=run_dir,
+                prefill_policy="round_robin",
+                decode_policy=_decode_policy_for(config.policy_name),
+                timeout_s=config.timeout_s,
+                router_request_timeout_s=(
+                    config.request_timeout_s
+                    if config.request_timeout_s is not None
+                    else config.timeout_s
+                ),
+                include_router=(
+                    config.mechanism_name in _mechanisms_with_router
+                ),
+                naive_dp=_naive_dp,
+            )
+            router_url = (
+                stack.router_url
+                if config.mechanism_name in _mechanisms_with_router
+                else None
+            )
+        else:
+            router_url = (
+                topology.router_url
+                if config.mechanism_name in _mechanisms_with_router
+                else None
+            )
+
+        metrics_path = run_dir / "request-metrics.jsonl"
+        replay_config = ReplayConfig(
+            trace_path=sampled_trace_path,
+            output_path=metrics_path,
+            policy_name=config.policy_name,
+            mechanism_name=config.mechanism_name,
+            topology=topology,
+            router_url=router_url,
+            model_name=topology.model_name,
+            pace=True,
+            time_scale=config.time_scale,
+            request_limit=None,
+            concurrency_limit=config.concurrency_limit,
+            header_mode=_header_mode_for(config.policy_name),
+            timeout_s=config.timeout_s,
+            stream=config.stream,
+            stream_idle_timeout_s=config.stream_idle_timeout_s,
+            kvcache_direct_max_uncached_tokens=config.kvcache_direct_max_uncached_tokens,
+            kvcache_admission_mode=config.kvcache_admission_mode,  # type: ignore[arg-type]
+            kvcache_seed_max_resident_tokens=config.kvcache_seed_max_resident_tokens,
+            kvcache_seed_max_output_tokens=config.kvcache_seed_max_output_tokens,
+            kvcache_seed_min_turn_id=config.kvcache_seed_min_turn_id,
+            kvcache_seed_only_multiturn_sessions=(
+                config.kvcache_seed_only_multiturn_sessions
+            ),
+            kvcache_prefill_backup_policy=config.kvcache_prefill_backup_policy,  # type: ignore[arg-type]
+            kvcache_seed_max_inflight_decode=(
+                config.kvcache_seed_max_inflight_decode
+            ),
+            kvcache_seed_max_decode_transfer_queue_reqs=(
+                config.kvcache_seed_max_decode_transfer_queue_reqs
+            ),
+            kvcache_direct_max_decode_transfer_queue_reqs=(
+                config.kvcache_direct_max_decode_transfer_queue_reqs
+            ),
+            kvcache_prefill_priority_eviction=(
+                config.kvcache_prefill_priority_eviction
+            ),
+            kvcache_prefill_direct_priority=config.kvcache_prefill_direct_priority,
+            kvcache_prefill_normal_priority=config.kvcache_prefill_normal_priority,
+            pool_poll_interval_s=config.pool_poll_interval_s,
+            pool_poll_include_sessions=config.pool_poll_include_sessions,
+            enable_backpressure=config.enable_backpressure,
+            backpressure_max_pause_s=config.backpressure_max_pause_s,
+            kvcache_migration_reject_threshold=config.kvcache_migration_reject_threshold,
+        )
+        if config.request_timeout_s is not None:
+            replay_config = replace(
+                replay_config,
+                timeout_s=config.request_timeout_s,
+            )
+        asyncio.run(replay_trace(replay_config))
+    finally:
+        signal.signal(signal.SIGINT, previous_sigint)
+        signal.signal(signal.SIGTERM, previous_sigterm)
+        if stack is not None:
+            stack.stop()
+
+    benchmark_config_path = run_dir / "benchmark-config.json"
+    with benchmark_config_path.open("w", encoding="utf-8") as handle:
+        json.dump(
+            {
+                "policy_name": config.policy_name,
+                "mechanism_name": config.mechanism_name,
+                "target_duration_s": config.target_duration_s,
+                "start_time_s": config.start_time_s,
+                "session_sample_rate": config.session_sample_rate,
+                "min_turns": config.min_turns,
+                "time_scale": config.time_scale,
+                "concurrency_limit": config.concurrency_limit,
+                "timeout_s": config.timeout_s,
+                "request_timeout_s": config.request_timeout_s,
+                "stream": config.stream,
+                "stream_idle_timeout_s": config.stream_idle_timeout_s,
+                "kvcache_direct_max_uncached_tokens": config.kvcache_direct_max_uncached_tokens,
+                "kvcache_admission_mode": config.kvcache_admission_mode,
+                "kvcache_seed_max_resident_tokens": config.kvcache_seed_max_resident_tokens,
+                "kvcache_seed_max_output_tokens": config.kvcache_seed_max_output_tokens,
+                "kvcache_seed_min_turn_id": config.kvcache_seed_min_turn_id,
+                "kvcache_seed_only_multiturn_sessions": (
+                    config.kvcache_seed_only_multiturn_sessions
+                ),
+                "kvcache_prefill_backup_policy": config.kvcache_prefill_backup_policy,
+                "kvcache_seed_max_inflight_decode": (
+                    config.kvcache_seed_max_inflight_decode
+                ),
+                "kvcache_seed_max_decode_transfer_queue_reqs": (
+                    config.kvcache_seed_max_decode_transfer_queue_reqs
+                ),
+                "kvcache_direct_max_decode_transfer_queue_reqs": (
+                    config.kvcache_direct_max_decode_transfer_queue_reqs
+                ),
+                "kvcache_prefill_priority_eviction": (
+                    config.kvcache_prefill_priority_eviction
+                ),
+                "kvcache_prefill_direct_priority": (
+                    config.kvcache_prefill_direct_priority
+                ),
+                "kvcache_prefill_normal_priority": (
+                    config.kvcache_prefill_normal_priority
+                ),
+                "pool_poll_interval_s": config.pool_poll_interval_s,
+                "pool_poll_include_sessions": config.pool_poll_include_sessions,
+                "enable_backpressure": config.enable_backpressure,
+                "backpressure_max_pause_s": config.backpressure_max_pause_s,
+                "kvcache_migration_reject_threshold": config.kvcache_migration_reject_threshold,
+                "sample_profile": config.sample_profile,
+                "min_initial_input_tokens": config.min_initial_input_tokens,
+                "max_initial_input_tokens": config.max_initial_input_tokens,
+                "max_append_input_tokens": config.max_append_input_tokens,
+                "max_output_tokens": config.max_output_tokens,
+                "min_overlap_ratio": config.min_overlap_ratio,
+                "sample_summary": asdict(sample_summary),
+                "topology": {
+                    "model_path": config.topology.model_path,
+                    "router_url": topology.router_url,
+                    "transfer_backend": topology.transfer_backend,
+                    "force_rdma": topology.force_rdma,
+                    "ib_device": topology.ib_device,
+                    "prefill_workers": [
+                        worker.worker_id for worker in topology.prefill_workers
+                    ],
+                    "decode_workers": [
+                        worker.worker_id for worker in topology.decode_workers
+                    ],
+                    "direct_workers": [
+                        worker.worker_id for worker in topology.direct_workers
+                    ],
+                },
+            },
+            handle,
+            indent=2,
+            sort_keys=True,
+        )
+
+    return BenchmarkArtifacts(
+        run_dir=run_dir,
+        sampled_trace_path=sampled_trace_path,
+        metrics_path=run_dir / "request-metrics.jsonl",
+        summary_path=run_dir / "request-metrics.jsonl.summary.json",
+        benchmark_config_path=benchmark_config_path,
+    )
+
+
+def _decode_policy_for(policy_name: str) -> str:
+    if policy_name == "sticky":
+        return "manual"
+    if policy_name == "kv-aware":
+        return "consistent_hashing"
+    return "round_robin"
+
+
+def _header_mode_for(policy_name: str) -> str:
+    if policy_name == "sticky":
+        return "routing-key"
+    if policy_name == "kv-aware":
+        return "target-worker"
+    return "none"

src/agentic_pd_hybrid/cli.py

@@ -0,0 +1,856 @@
+from __future__ import annotations
+
+import argparse
+import asyncio
+from pathlib import Path
+
+from agentic_pd_hybrid.benchmark import BenchmarkConfig, run_live_benchmark
+from agentic_pd_hybrid.launcher import build_launch_plan
+from agentic_pd_hybrid.microbench import SmallAppendTraceConfig, write_small_append_trace
+from agentic_pd_hybrid.profile import ProfileConfig, print_profile_summary, write_profile
+from agentic_pd_hybrid.replay import ReplayConfig, replay_trace
+from agentic_pd_hybrid.sampling import SessionSampleConfig, sample_trace_sessions
+from agentic_pd_hybrid.trace_profiles import (
+    NormalizeTraceLengthsConfig,
+    normalize_trace_lengths,
+)
+from agentic_pd_hybrid.topology import build_single_node_topology
+
+
+def _normalize_mechanism_name(name: str) -> str:
+    normalized = name.strip().lower()
+    aliases = {
+        "pd-disagg": "pd-disaggregation",
+        "pd-disaggregation": "pd-disaggregation",
+        "pd-hybrid": "pd-disaggregation",
+        "baseline-pd-disagg": "pd-disaggregation",
+        "pd-colo": "pd-colo",
+        "direct-d": "pd-colo",
+        "colocation": "pd-colo",
+        "kvcache-centric": "kvcache-centric",
+        "turn2+-direct-to-d": "kvcache-centric",
+        "pd-with-d-append": "kvcache-centric",
+    }
+    if normalized not in aliases:
+        raise ValueError(f"Unsupported mechanism: {name}")
+    return aliases[normalized]
+
+
+def _parse_gpu_id_list(value: str | None) -> tuple[int, ...] | None:
+    if value is None:
+        return None
+    items = [item.strip() for item in value.split(",") if item.strip()]
+    if not items:
+        return tuple()
+    return tuple(int(item) for item in items)
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser(description="Agentic PD hybrid prototype")
+    subparsers = parser.add_subparsers(dest="command", required=True)
+
+    print_launch = subparsers.add_parser(
+        "print-launch",
+        help="Print one-node SGLang PD launch commands",
+    )
+    _add_topology_arguments(print_launch)
+    print_launch.add_argument("--prefill-policy", default="round_robin")
+    print_launch.add_argument("--decode-policy", default="manual")
+
+    replay = subparsers.add_parser(
+        "replay",
+        help="Replay trace and log request-level metrics",
+    )
+    _add_topology_arguments(replay)
+    replay.add_argument("--trace", type=Path, required=True)
+    replay.add_argument("--output", type=Path, required=True)
+    replay.add_argument(
+        "--policy",
+        choices=["default", "sticky", "kv-aware"],
+        default="sticky",
+    )
+    replay.add_argument(
+        "--mechanism",
+        choices=[
+            "pd-disaggregation",
+            "pd-hybrid",
+            "pd-disagg",
+            "pd-colo",
+            "direct-d",
+            "kvcache-centric",
+            "turn2+-direct-to-d",
+            "pd-with-d-append",
+        ],
+        default="pd-disaggregation",
+    )
+    replay.add_argument("--router-url")
+    replay.add_argument("--model")
+    replay.add_argument(
+        "--header-mode",
+        choices=["auto", "none", "routing-key", "target-worker"],
+        default="auto",
+    )
+    replay.add_argument(
+        "--request-limit",
+        type=int,
+        default=None,
+        help="Replay at most this many requests",
+    )
+    replay.add_argument(
+        "--no-pace",
+        action="store_true",
+        help="Disable wall-clock pacing from trace timestamps",
+    )
+    replay.add_argument(
+        "--time-scale",
+        type=float,
+        default=1.0,
+        help="Scale trace timing by this factor when pacing is enabled",
+    )
+    replay.add_argument(
+        "--concurrency-limit",
+        type=int,
+        default=32,
+    )
+    replay.add_argument(
+        "--timeout-s",
+        type=float,
+        default=600.0,
+    )
+    replay.add_argument(
+        "--no-stream",
+        action="store_true",
+        help="Use non-streaming OpenAI responses for more robust E2E-only replay.",
+    )
+    replay.add_argument(
+        "--stream-idle-timeout-s",
+        type=float,
+        default=900.0,
+        help="Abort a streaming request if no SSE line arrives within this many seconds.",
+    )
+    replay.add_argument(
+        "--kvcache-direct-max-uncached-tokens",
+        type=int,
+        default=2048,
+        help="For kvcache-centric routing, bypass P when the uncached suffix is at most this many tokens.",
+    )
+    replay.add_argument(
+        "--kvcache-admission-mode",
+        choices=["router", "worker"],
+        default="router",
+        help=(
+            "For kvcache-centric routing, use router shadow-state admission "
+            "or query the decode worker on the critical path."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-seed-max-resident-tokens",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric routing, do not seed/reseed a decode session "
+            "when input+output tokens exceed this value."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-seed-max-output-tokens",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric routing, do not seed/reseed a decode session "
+            "when output tokens exceed this value."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-seed-min-turn-id",
+        type=int,
+        default=1,
+        help=(
+            "For kvcache-centric routing, do not seed/reseed a decode session "
+            "before this turn id."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-seed-only-multiturn-sessions",
+        action="store_true",
+        help=(
+            "Oracle ablation for kvcache-centric routing: only seed sessions "
+            "that have more than one turn in the replay trace."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-prefill-backup-policy",
+        choices=["release-after-transfer", "capacity-backup"],
+        default="release-after-transfer",
+        help=(
+            "For kvcache-centric seed/reseed, release the P-side session after "
+            "P->D transfer or keep a capacity-limited P-side backup."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-seed-max-inflight-decode",
+        type=int,
+        default=3,
+        help=(
+            "For kvcache-centric routing, skip seed/reseed when the router "
+            "shadow inflight decode load assigned to the target D exceeds this value. "
+            "Use a negative value to disable this filter."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-seed-max-decode-transfer-queue-reqs",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric worker admission, skip seed/reseed when the "
+            "target D reports more transfer-queue requests than this value. "
+            "Use a negative value to disable this filter."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-direct-max-decode-transfer-queue-reqs",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric worker admission, skip direct-to-D append when "
+            "the target D reports more transfer-queue requests than this value. "
+            "Use a negative value to disable this filter."
+        ),
+    )
+    replay.add_argument(
+        "--kvcache-prefill-priority-eviction",
+        action="store_true",
+        help=(
+            "For kvcache-centric routing, mark P-side prefixes predicted to move "
+            "direct-to-D as lower priority. Requires P workers to use "
+            "--radix-eviction-policy priority."
+        ),
+    )
+    replay.add_argument("--kvcache-prefill-direct-priority", type=int, default=-100)
+    replay.add_argument("--kvcache-prefill-normal-priority", type=int, default=100)
+    replay.add_argument(
+        "--pool-poll-interval-s",
+        type=float,
+        default=0.0,
+        help=(
+            "Poll each P/D worker's /server_info every N seconds and write a "
+            "time-series snapshot to <run_dir>/d-pool-timeseries.jsonl. "
+            "0 disables polling."
+        ),
+    )
+    replay.add_argument(
+        "--pool-poll-no-sessions",
+        action="store_true",
+        help=(
+            "Disable per-session detail in the pool timeseries (smaller files)."
+        ),
+    )
+    replay.add_argument(
+        "--enable-backpressure",
+        action="store_true",
+        help=(
+            "Honor recommended_pause_ms hints from D's admission endpoint. "
+            "When set, replay sleeps before issuing requests to a saturated D. "
+            "Default off — preserves baseline behavior."
+        ),
+    )
+    replay.add_argument(
+        "--backpressure-max-pause-s",
+        type=float,
+        default=2.0,
+        help="Cap on per-request backpressure sleep, regardless of D hint.",
+    )
+    replay.add_argument(
+        "--kvcache-migration-reject-threshold",
+        type=int,
+        default=3,
+        help=(
+            "Per-(session, D) admission-reject count after which KvAwarePolicy "
+            "skips that D for the session (forces migration). 0 disables. "
+            "See REFACTOR_PLAN_V1 §6.2 / TEAM_REPORT §2.1."
+        ),
+    )
+
+    sample = subparsers.add_parser(
+        "sample-sessions",
+        help="Sample a session-granularity trace shard for live benchmarking",
+    )
+    sample.add_argument("--trace", type=Path, required=True)
+    sample.add_argument("--output", type=Path, required=True)
+    sample.add_argument("--target-duration-s", type=float, default=600.0)
+    sample.add_argument("--start-time-s", type=float, default=0.0)
+    sample.add_argument("--session-sample-rate", type=float, default=0.01)
+    sample.add_argument("--min-turns", type=int, default=1)
+    sample.add_argument("--max-requests", type=int, default=None)
+    sample.add_argument(
+        "--profile",
+        choices=["default", "small-append"],
+        default="default",
+        help="Optional workload-shape filter for live benchmarks.",
+    )
+    sample.add_argument("--min-initial-input-tokens", type=int, default=None)
+    sample.add_argument("--max-initial-input-tokens", type=int, default=None)
+    sample.add_argument("--max-append-input-tokens", type=int, default=None)
+    sample.add_argument("--max-output-tokens", type=int, default=None)
+    sample.add_argument("--min-overlap-ratio", type=float, default=None)
+
+    normalize = subparsers.add_parser(
+        "normalize-trace-lengths",
+        help="Rewrite a trace to a fixed turn1/append/output length profile",
+    )
+    normalize.add_argument("--trace", type=Path, required=True)
+    normalize.add_argument("--output", type=Path, required=True)
+    normalize.add_argument("--initial-input-length", type=int, default=10_000)
+    normalize.add_argument("--append-input-length", type=int, default=1_000)
+    normalize.add_argument("--output-length", type=int, default=1_000)
+    normalize.add_argument("--max-requests", type=int, default=None)
+
+    profile = subparsers.add_parser(
+        "profile",
+        help="Profile a trace and optional request metrics for routing analysis",
+    )
+    profile.add_argument("--trace", type=Path, required=True)
+    profile.add_argument("--output", type=Path, default=None)
+    profile.add_argument("--metrics", type=Path, default=None)
+    profile.add_argument("--baseline-metrics", type=Path, default=None)
+    profile.add_argument("--candidate-metrics", type=Path, default=None)
+    profile.add_argument("--direct-max-uncached-tokens", type=int, default=2048)
+
+    micro = subparsers.add_parser(
+        "make-small-append-trace",
+        help="Generate a synthetic multi-turn trace with small turn2+ appends",
+    )
+    micro.add_argument("--output", type=Path, required=True)
+    micro.add_argument("--session-count", type=int, default=8)
+    micro.add_argument("--turns-per-session", type=int, default=3)
+    micro.add_argument("--initial-input-length", type=int, default=10_000)
+    micro.add_argument("--append-input-length", type=int, default=1_000)
+    micro.add_argument("--output-length", type=int, default=1_000)
+    micro.add_argument("--inter-turn-gap-s", type=float, default=1.0)
+    micro.add_argument("--session-stagger-s", type=float, default=0.1)
+
+    benchmark = subparsers.add_parser(
+        "benchmark-live",
+        help="Launch a real PD stack, sample sessions, and collect live E2E numbers",
+    )
+    _add_topology_arguments(benchmark)
+    benchmark.add_argument("--trace", type=Path, required=True)
+    benchmark.add_argument(
+        "--policy",
+        choices=["default", "sticky", "kv-aware"],
+        default="sticky",
+    )
+    benchmark.add_argument(
+        "--mechanism",
+        choices=[
+            "pd-disaggregation",
+            "pd-hybrid",
+            "pd-disagg",
+            "pd-colo",
+            "direct-d",
+            "kvcache-centric",
+            "turn2+-direct-to-d",
+            "pd-with-d-append",
+        ],
+        default="pd-disaggregation",
+    )
+    benchmark.add_argument("--output-root", type=Path, default=Path("outputs/live"))
+    benchmark.add_argument("--target-duration-s", type=float, default=600.0)
+    benchmark.add_argument("--start-time-s", type=float, default=0.0)
+    benchmark.add_argument("--session-sample-rate", type=float, default=0.01)
+    benchmark.add_argument("--min-turns", type=int, default=1)
+    benchmark.add_argument("--time-scale", type=float, default=1.0)
+    benchmark.add_argument("--concurrency-limit", type=int, default=32)
+    benchmark.add_argument("--timeout-s", type=float, default=1200.0)
+    benchmark.add_argument(
+        "--request-timeout-s",
+        type=float,
+        default=None,
+        help=(
+            "Per-request replay/router timeout. If unset, --timeout-s is used. "
+            "--timeout-s still controls stack startup."
+        ),
+    )
+    benchmark.add_argument(
+        "--no-stream",
+        action="store_true",
+        help="Use non-streaming OpenAI responses for E2E-only live benchmarking.",
+    )
+    benchmark.add_argument(
+        "--stream-idle-timeout-s",
+        type=float,
+        default=900.0,
+        help="Abort a streaming request if no SSE line arrives within this many seconds.",
+    )
+    benchmark.add_argument(
+        "--kvcache-direct-max-uncached-tokens",
+        type=int,
+        default=2048,
+        help="For kvcache-centric routing, bypass P when the uncached suffix is at most this many tokens.",
+    )
+    benchmark.add_argument(
+        "--kvcache-admission-mode",
+        choices=["router", "worker"],
+        default="router",
+        help=(
+            "For kvcache-centric routing, use router shadow-state admission "
+            "or query the decode worker on the critical path."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-seed-max-resident-tokens",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric routing, do not seed/reseed a decode session "
+            "when input+output tokens exceed this value."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-seed-max-output-tokens",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric routing, do not seed/reseed a decode session "
+            "when output tokens exceed this value."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-seed-min-turn-id",
+        type=int,
+        default=1,
+        help=(
+            "For kvcache-centric routing, do not seed/reseed a decode session "
+            "before this turn id."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-seed-only-multiturn-sessions",
+        action="store_true",
+        help=(
+            "Oracle ablation for kvcache-centric routing: only seed sessions "
+            "that have more than one turn in the replay trace."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-prefill-backup-policy",
+        choices=["release-after-transfer", "capacity-backup"],
+        default="release-after-transfer",
+        help=(
+            "For kvcache-centric seed/reseed, release the P-side session after "
+            "P->D transfer or keep a capacity-limited P-side backup."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-seed-max-inflight-decode",
+        type=int,
+        default=3,
+        help=(
+            "For kvcache-centric routing, skip seed/reseed when the router "
+            "shadow inflight decode load assigned to the target D exceeds this value. "
+            "Use a negative value to disable this filter."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-seed-max-decode-transfer-queue-reqs",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric worker admission, skip seed/reseed when the "
+            "target D reports more transfer-queue requests than this value. "
+            "Use a negative value to disable this filter."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-direct-max-decode-transfer-queue-reqs",
+        type=int,
+        default=None,
+        help=(
+            "For kvcache-centric worker admission, skip direct-to-D append when "
+            "the target D reports more transfer-queue requests than this value. "
+            "Use a negative value to disable this filter."
+        ),
+    )
+    benchmark.add_argument(
+        "--kvcache-prefill-priority-eviction",
+        action="store_true",
+        help=(
+            "For kvcache-centric benchmark-live, launch P workers with priority "
+            "radix eviction and mark direct-to-D predicted prefixes lower priority."
+        ),
+    )
+    benchmark.add_argument("--kvcache-prefill-direct-priority", type=int, default=-100)
+    benchmark.add_argument("--kvcache-prefill-normal-priority", type=int, default=100)
+    benchmark.add_argument(
+        "--pool-poll-interval-s",
+        type=float,
+        default=0.0,
+        help=(
+            "Poll each P/D worker's /server_info every N seconds and write a "
+            "time-series snapshot to <run_dir>/d-pool-timeseries.jsonl. "
+            "0 disables polling."
+        ),
+    )
+    benchmark.add_argument(
+        "--pool-poll-no-sessions",
+        action="store_true",
+        help=(
+            "Disable per-session detail in the pool timeseries (smaller files)."
+        ),
+    )
+    benchmark.add_argument(
+        "--enable-backpressure",
+        action="store_true",
+        help=(
+            "Honor recommended_pause_ms hints from D's admission endpoint."
+        ),
+    )
+    benchmark.add_argument(
+        "--backpressure-max-pause-s",
+        type=float,
+        default=2.0,
+        help="Cap on per-request backpressure sleep, regardless of D hint.",
+    )
+    benchmark.add_argument(
+        "--kvcache-migration-reject-threshold",
+        type=int,
+        default=3,
+        help=(
+            "Per-(session, D) admission-reject count after which KvAwarePolicy "
+            "skips that D for the session (forces migration). 0 disables. "
+            "See REFACTOR_PLAN_V1 §6.2 / TEAM_REPORT §2.1."
+        ),
+    )
+    benchmark.add_argument(
+        "--sample-profile",
+        choices=["default", "small-append"],
+        default="default",
+        help="Optional session-shape filter applied before live replay.",
+    )
+    benchmark.add_argument("--min-initial-input-tokens", type=int, default=None)
+    benchmark.add_argument("--max-initial-input-tokens", type=int, default=None)
+    benchmark.add_argument("--max-append-input-tokens", type=int, default=None)
+    benchmark.add_argument("--max-output-tokens", type=int, default=None)
+    benchmark.add_argument("--min-overlap-ratio", type=float, default=None)
+
+    args = parser.parse_args()
+
+    if args.command == "print-launch":
+        topology = _topology_from_args(args)
+        has_pd = bool(topology.prefill_workers and topology.decode_workers)
+        has_direct_only = bool(
+            topology.direct_workers
+            and not topology.prefill_workers
+            and not topology.decode_workers
+        )
+        plan = build_launch_plan(
+            topology,
+            prefill_policy=args.prefill_policy,
+            decode_policy=args.decode_policy,
+            include_router=has_pd or has_direct_only,
+            naive_dp=has_direct_only,
+        )
+        print(plan.render())
+        return
+
+    if args.command == "replay":
+        topology = _topology_from_args(args)
+        config = ReplayConfig(
+            trace_path=args.trace,
+            output_path=args.output,
+            policy_name=args.policy,
+            mechanism_name=_normalize_mechanism_name(args.mechanism),
+            topology=topology,
+            router_url=args.router_url,
+            model_name=args.model,
+            pace=not args.no_pace,
+            time_scale=args.time_scale,
+            request_limit=args.request_limit,
+            concurrency_limit=args.concurrency_limit,
+            header_mode=args.header_mode,
+            timeout_s=args.timeout_s,
+            stream=not args.no_stream,
+            stream_idle_timeout_s=args.stream_idle_timeout_s,
+            kvcache_direct_max_uncached_tokens=args.kvcache_direct_max_uncached_tokens,
+            kvcache_admission_mode=args.kvcache_admission_mode,
+            kvcache_seed_max_resident_tokens=args.kvcache_seed_max_resident_tokens,
+            kvcache_seed_max_output_tokens=args.kvcache_seed_max_output_tokens,
+            kvcache_seed_min_turn_id=args.kvcache_seed_min_turn_id,
+            kvcache_seed_only_multiturn_sessions=(
+                args.kvcache_seed_only_multiturn_sessions
+            ),
+            kvcache_prefill_backup_policy=args.kvcache_prefill_backup_policy,
+            kvcache_seed_max_inflight_decode=(
+                None
+                if args.kvcache_seed_max_inflight_decode < 0
+                else args.kvcache_seed_max_inflight_decode
+            ),
+            kvcache_seed_max_decode_transfer_queue_reqs=(
+                None
+                if args.kvcache_seed_max_decode_transfer_queue_reqs is None
+                or args.kvcache_seed_max_decode_transfer_queue_reqs < 0
+                else args.kvcache_seed_max_decode_transfer_queue_reqs
+            ),
+            kvcache_direct_max_decode_transfer_queue_reqs=(
+                None
+                if args.kvcache_direct_max_decode_transfer_queue_reqs is None
+                or args.kvcache_direct_max_decode_transfer_queue_reqs < 0
+                else args.kvcache_direct_max_decode_transfer_queue_reqs
+            ),
+            kvcache_prefill_priority_eviction=(
+                args.kvcache_prefill_priority_eviction
+            ),
+            kvcache_prefill_direct_priority=args.kvcache_prefill_direct_priority,
+            kvcache_prefill_normal_priority=args.kvcache_prefill_normal_priority,
+            pool_poll_interval_s=args.pool_poll_interval_s,
+            pool_poll_include_sessions=not args.pool_poll_no_sessions,
+            enable_backpressure=args.enable_backpressure,
+            backpressure_max_pause_s=args.backpressure_max_pause_s,
+            kvcache_migration_reject_threshold=args.kvcache_migration_reject_threshold,
+        )
+        results = asyncio.run(replay_trace(config))
+        print(
+            f"wrote {len(results)} request records to {args.output} and "
+            f"{args.output}{'.summary.json'}"
+        )
+        return
+
+    if args.command == "profile":
+        if (args.baseline_metrics is None) != (args.candidate_metrics is None):
+            raise ValueError(
+                "--baseline-metrics and --candidate-metrics must be provided together"
+            )
+        report = write_profile(
+            ProfileConfig(
+                trace_path=args.trace,
+                output_path=args.output,
+                metrics_path=args.metrics,
+                baseline_metrics_path=args.baseline_metrics,
+                candidate_metrics_path=args.candidate_metrics,
+                direct_max_uncached_tokens=args.direct_max_uncached_tokens,
+            )
+        )
+        print_profile_summary(report)
+        if args.output is not None:
+            print(f"wrote profile to {args.output}")
+        return
+
+    if args.command == "sample-sessions":
+        summary = sample_trace_sessions(
+            SessionSampleConfig(
+                trace_path=args.trace,
+                output_path=args.output,
+                target_duration_s=args.target_duration_s,
+                start_time_s=args.start_time_s,
+                session_sample_rate=args.session_sample_rate,
+                min_turns=args.min_turns,
+                max_requests=args.max_requests,
+                profile=args.profile,
+                min_initial_input_tokens=args.min_initial_input_tokens,
+                max_initial_input_tokens=args.max_initial_input_tokens,
+                max_append_input_tokens=args.max_append_input_tokens,
+                max_output_tokens=args.max_output_tokens,
+                min_overlap_ratio=args.min_overlap_ratio,
+            )
+        )
+        print(
+            f"wrote {summary.request_count} requests from {summary.session_count} sessions "
+            f"covering {summary.sampled_duration_s:.3f}s to {args.output}"
+        )
+        return
+
+    if args.command == "normalize-trace-lengths":
+        summary = normalize_trace_lengths(
+            NormalizeTraceLengthsConfig(
+                trace_path=args.trace,
+                output_path=args.output,
+                initial_input_length=args.initial_input_length,
+                append_input_length=args.append_input_length,
+                output_length=args.output_length,
+                max_requests=args.max_requests,
+            )
+        )
+        print(
+            f"wrote {summary.request_count} normalized requests from "
+            f"{summary.session_count} sessions to {args.output}"
+        )
+        return
+
+    if args.command == "make-small-append-trace":
+        summary = write_small_append_trace(
+            SmallAppendTraceConfig(
+                output_path=args.output,
+                session_count=args.session_count,
+                turns_per_session=args.turns_per_session,
+                initial_input_length=args.initial_input_length,
+                append_input_length=args.append_input_length,
+                output_length=args.output_length,
+                inter_turn_gap_s=args.inter_turn_gap_s,
+                session_stagger_s=args.session_stagger_s,
+            )
+        )
+        print(
+            f"wrote {summary.request_count} requests across {summary.session_count} sessions "
+            f"to {args.output}"
+        )
+        return
+
+    if args.command == "benchmark-live":
+        topology = _topology_from_args(args)
+        artifacts = run_live_benchmark(
+            BenchmarkConfig(
+                trace_path=args.trace,
+                output_root=args.output_root,
+                topology=topology,
+                policy_name=args.policy,
+                mechanism_name=_normalize_mechanism_name(args.mechanism),
+                target_duration_s=args.target_duration_s,
+                start_time_s=args.start_time_s,
+                session_sample_rate=args.session_sample_rate,
+                min_turns=args.min_turns,
+                time_scale=args.time_scale,
+                concurrency_limit=args.concurrency_limit,
+                timeout_s=args.timeout_s,
+                request_timeout_s=args.request_timeout_s,
+                stream=not args.no_stream,
+                stream_idle_timeout_s=args.stream_idle_timeout_s,
+                kvcache_direct_max_uncached_tokens=args.kvcache_direct_max_uncached_tokens,
+                kvcache_admission_mode=args.kvcache_admission_mode,
+                kvcache_seed_max_resident_tokens=args.kvcache_seed_max_resident_tokens,
+                kvcache_seed_max_output_tokens=args.kvcache_seed_max_output_tokens,
+                kvcache_seed_min_turn_id=args.kvcache_seed_min_turn_id,
+                kvcache_seed_only_multiturn_sessions=(
+                    args.kvcache_seed_only_multiturn_sessions
+                ),
+                kvcache_prefill_backup_policy=args.kvcache_prefill_backup_policy,
+                kvcache_seed_max_inflight_decode=(
+                    None
+                    if args.kvcache_seed_max_inflight_decode < 0
+                    else args.kvcache_seed_max_inflight_decode
+                ),
+                kvcache_seed_max_decode_transfer_queue_reqs=(
+                    None
+                    if args.kvcache_seed_max_decode_transfer_queue_reqs is None
+                    or args.kvcache_seed_max_decode_transfer_queue_reqs < 0
+                    else args.kvcache_seed_max_decode_transfer_queue_reqs
+                ),
+                kvcache_direct_max_decode_transfer_queue_reqs=(
+                    None
+                    if args.kvcache_direct_max_decode_transfer_queue_reqs is None
+                    or args.kvcache_direct_max_decode_transfer_queue_reqs < 0
+                    else args.kvcache_direct_max_decode_transfer_queue_reqs
+                ),
+                kvcache_prefill_priority_eviction=(
+                    args.kvcache_prefill_priority_eviction
+                ),
+                kvcache_prefill_direct_priority=(
+                    args.kvcache_prefill_direct_priority
+                ),
+                kvcache_prefill_normal_priority=(
+                    args.kvcache_prefill_normal_priority
+                ),
+                pool_poll_interval_s=args.pool_poll_interval_s,
+                pool_poll_include_sessions=not args.pool_poll_no_sessions,
+                enable_backpressure=args.enable_backpressure,
+                backpressure_max_pause_s=args.backpressure_max_pause_s,
+                kvcache_migration_reject_threshold=args.kvcache_migration_reject_threshold,
+                sample_profile=args.sample_profile,
+                min_initial_input_tokens=args.min_initial_input_tokens,
+                max_initial_input_tokens=args.max_initial_input_tokens,
+                max_append_input_tokens=args.max_append_input_tokens,
+                max_output_tokens=args.max_output_tokens,
+                min_overlap_ratio=args.min_overlap_ratio,
+                launch_stack=True,
+            )
+        )
+        print(
+            f"benchmark artifacts written under {artifacts.run_dir}; "
+            f"metrics={artifacts.metrics_path} summary={artifacts.summary_path}"
+        )
+        return
+
+    raise AssertionError(f"Unhandled command: {args.command}")
+
+
+def _add_topology_arguments(parser: argparse.ArgumentParser) -> None:
+    parser.add_argument(
+        "--model-path",
+        default="~/models/Qwen/Qwen3-Coder-30B-A3B-Instruct",
+    )
+    parser.add_argument("--prefill-workers", type=int, default=1)
+    parser.add_argument("--decode-workers", type=int, default=1)
+    parser.add_argument("--direct-workers", type=int, default=0)
+    parser.add_argument("--prefill-tp-size", type=int, default=1)
+    parser.add_argument("--decode-tp-size", type=int, default=1)
+    parser.add_argument("--direct-tp-size", type=int, default=1)
+    parser.add_argument("--gpu-budget", type=int, default=8)
+    parser.add_argument(
+        "--prefill-gpu-ids",
+        default=None,
+        help="Comma-separated GPU IDs for prefill workers, e.g. 3,4",
+    )
+    parser.add_argument(
+        "--decode-gpu-ids",
+        default=None,
+        help="Comma-separated GPU IDs for decode workers, e.g. 5",
+    )
+    parser.add_argument(
+        "--direct-gpu-ids",
+        default=None,
+        help="Comma-separated GPU IDs for direct workers, e.g. 6",
+    )
+    parser.add_argument("--host", default="127.0.0.1")
+    parser.add_argument("--router-port", type=int, default=8000)
+    parser.add_argument("--prefill-port-base", type=int, default=30000)
+    parser.add_argument("--decode-port-base", type=int, default=31000)
+    parser.add_argument("--direct-port-base", type=int, default=32000)
+    parser.add_argument("--bootstrap-port-base", type=int, default=8998)
+    parser.add_argument("--transfer-backend", default="nixl")
+    parser.add_argument(
+        "--force-rdma",
+        action="store_true",
+        help=(
+            "Force real RDMA transport for PD KV transfer. "
+            "Currently this requires Mooncake plus --ib-device."
+        ),
+    )
+    parser.add_argument("--ib-device", default=None)
+    parser.add_argument(
+        "--no-trust-remote-code",
+        action="store_true",
+    )
+
+
+def _topology_from_args(args: argparse.Namespace):
+    transfer_backend = args.transfer_backend
+    if args.force_rdma:
+        transfer_backend = "mooncake"
+
+    return build_single_node_topology(
+        model_path=str(Path(args.model_path).expanduser()),
+        prefill_worker_count=args.prefill_workers,
+        decode_worker_count=args.decode_workers,
+        direct_worker_count=args.direct_workers,
+        prefill_tp_size=args.prefill_tp_size,
+        decode_tp_size=args.decode_tp_size,
+        direct_tp_size=args.direct_tp_size,
+        prefill_gpu_ids=_parse_gpu_id_list(args.prefill_gpu_ids),
+        decode_gpu_ids=_parse_gpu_id_list(args.decode_gpu_ids),
+        direct_gpu_ids=_parse_gpu_id_list(args.direct_gpu_ids),
+        total_gpu_budget=args.gpu_budget,
+        host=args.host,
+        router_port=args.router_port,
+        prefill_port_base=args.prefill_port_base,
+        decode_port_base=args.decode_port_base,
+        direct_port_base=args.direct_port_base,
+        bootstrap_port_base=args.bootstrap_port_base,
+        transfer_backend=transfer_backend,
+        force_rdma=args.force_rdma,
+        trust_remote_code=not args.no_trust_remote_code,
+        ib_device=args.ib_device,
+        direct_extra_server_args=("--enable-streaming-session",),
+    )
+
+
+if __name__ == "__main__":
+    main()

src/agentic_pd_hybrid/launcher.py

@@ -0,0 +1,185 @@
+from __future__ import annotations
+
+import shlex
+import sys
+from dataclasses import dataclass
+
+from agentic_pd_hybrid.topology import SingleNodeTopology, WorkerSpec
+
+
+@dataclass(frozen=True)
+class LaunchPlan:
+    prefill_commands: tuple[tuple[str, ...], ...]
+    decode_commands: tuple[tuple[str, ...], ...]
+    direct_commands: tuple[tuple[str, ...], ...]
+    router_command: tuple[str, ...] | None
+
+    def render(self) -> str:
+        sections: list[str] = []
+        for idx, command in enumerate(self.prefill_commands):
+            sections.append(_render_named_command(f"prefill-{idx}", command))
+        for idx, command in enumerate(self.decode_commands):
+            sections.append(_render_named_command(f"decode-{idx}", command))
+        for idx, command in enumerate(self.direct_commands):
+            sections.append(_render_named_command(f"direct-{idx}", command))
+        if self.router_command is not None:
+            sections.append(_render_named_command("router", self.router_command))
+        return "\n\n".join(sections)
+
+
+def build_launch_plan(
+    topology: SingleNodeTopology,
+    *,
+    prefill_policy: str = "round_robin",
+    decode_policy: str = "manual",
+    include_router: bool = True,
+    router_request_timeout_s: float | None = None,
+    naive_dp: bool = False,
+) -> LaunchPlan:
+    router_command: tuple[str, ...] | None = None
+    if include_router:
+        if topology.prefill_workers and topology.decode_workers:
+            router_command = _build_router_command(
+                topology,
+                prefill_policy=prefill_policy,
+                decode_policy=decode_policy,
+                request_timeout_s=router_request_timeout_s,
+            )
+        elif naive_dp and topology.direct_workers:
+            router_command = _build_dp_router_command(
+                topology,
+                backend_policy=decode_policy,
+                request_timeout_s=router_request_timeout_s,
+            )
+
+    return LaunchPlan(
+        prefill_commands=tuple(
+            _build_server_command(topology, worker) for worker in topology.prefill_workers
+        ),
+        decode_commands=tuple(
+            _build_server_command(topology, worker) for worker in topology.decode_workers
+        ),
+        direct_commands=tuple(
+            _build_server_command(topology, worker, naive_dp=naive_dp)
+            for worker in topology.direct_workers
+        ),
+        router_command=router_command,
+    )
+
+
+def _build_server_command(
+    topology: SingleNodeTopology,
+    worker: WorkerSpec,
+    naive_dp: bool = False,
+) -> tuple[str, ...]:
+    command = [
+        sys.executable,
+        "-B",
+        "-u",
+        "-m",
+        "sglang.launch_server",
+        "--model-path",
+        topology.model_path,
+        "--host",
+        worker.host,
+        "--port",
+        str(worker.port),
+        "--base-gpu-id",
+        str(worker.gpu_id),
+    ]
+    # Naive DP direct workers: no disaggregation flags at all
+    if not (naive_dp and worker.role == "direct"):
+        command.extend([
+            "--disaggregation-mode",
+            _disaggregation_mode_for(worker),
+            "--disaggregation-transfer-backend",
+            topology.transfer_backend,
+        ])
+    if worker.tp_size > 1:
+        command.extend(["--tp-size", str(worker.tp_size)])
+    if topology.trust_remote_code:
+        command.append("--trust-remote-code")
+    command.append("--enable-cache-report")
+    if worker.bootstrap_port is not None:
+        command.extend(
+            ["--disaggregation-bootstrap-port", str(worker.bootstrap_port)]
+        )
+    if topology.ib_device:
+        command.extend(["--disaggregation-ib-device", topology.ib_device])
+    command.extend(topology.extra_server_args)
+    if worker.role == "prefill":
+        command.extend(topology.prefill_extra_server_args)
+    elif worker.role == "decode":
+        command.extend(topology.decode_extra_server_args)
+    else:
+        command.extend(topology.direct_extra_server_args)
+    return tuple(command)
+
+
+def _build_router_command(
+    topology: SingleNodeTopology,
+    *,
+    prefill_policy: str,
+    decode_policy: str,
+    request_timeout_s: float | None,
+) -> tuple[str, ...]:
+    command: list[str] = [
+        sys.executable,
+        "-B",
+        "-u",
+        "-m",
+        "agentic_pd_hybrid.pd_router",
+        "--host",
+        topology.router_host,
+        "--port",
+        str(topology.router_port),
+        "--prefill-policy",
+        prefill_policy,
+        "--decode-policy",
+        decode_policy,
+    ]
+    if request_timeout_s is not None:
+        command.extend(["--request-timeout-s", str(request_timeout_s)])
+    for worker in topology.prefill_workers:
+        command.extend(
+            ["--prefill", worker.url, str(worker.bootstrap_port or topology.router_port)]
+        )
+    for worker in topology.decode_workers:
+        command.extend(["--decode", worker.url])
+    return tuple(command)
+
+
+def _build_dp_router_command(
+    topology: SingleNodeTopology,
+    *,
+    backend_policy: str,
+    request_timeout_s: float | None,
+) -> tuple[str, ...]:
+    command: list[str] = [
+        sys.executable,
+        "-B",
+        "-u",
+        "-m",
+        "agentic_pd_hybrid.pd_router",
+        "--host",
+        topology.router_host,
+        "--port",
+        str(topology.router_port),
+        "--backend-policy",
+        backend_policy,
+    ]
+    if request_timeout_s is not None:
+        command.extend(["--request-timeout-s", str(request_timeout_s)])
+    for worker in topology.direct_workers:
+        command.extend(["--backend", worker.url])
+    return tuple(command)
+
+
+def _render_named_command(name: str, command: tuple[str, ...]) -> str:
+    return f"# {name}\n" + " ".join(shlex.quote(part) for part in command)
+
+
+def _disaggregation_mode_for(worker: WorkerSpec) -> str:
+    if worker.role == "direct":
+        return "null"
+    return worker.role

src/agentic_pd_hybrid/metrics.py

@@ -0,0 +1,229 @@
+from __future__ import annotations
+
+import json
+import statistics
+from collections import Counter
+from dataclasses import asdict, dataclass
+from pathlib import Path
+from typing import Any
+
+from agentic_pd_hybrid.policies import RoutingDecision
+from agentic_pd_hybrid.trace import TraceRequest
+
+
+@dataclass(frozen=True)
+class RequestMetrics:
+    request_id: str
+    session_id: str
+    turn_id: int
+    mechanism_name: str
+    execution_mode: str
+    trace_timestamp_s: float
+    input_length: int
+    output_length: int
+    request_type: str
+    policy_name: str
+    assigned_prefill_node: str
+    assigned_decode_node: str
+    assigned_decode_index: int
+    inflight_decode_load_at_assignment: int
+    reuse_expected: bool
+    reuse_observed: bool
+    observed_overlap_blocks: int
+    kv_transfer_blocks: int
+    actual_kv_transfer_blocks: int
+    cached_tokens: int
+    prefill_request_priority: int | None
+    decode_request_priority: int | None
+    re_prefill_required: bool
+    effective_input_length: int | None
+    session_reused: bool
+    session_reset: bool
+    latency_s: float | None
+    ttft_s: float | None
+    tpot_s: float | None
+    error: str | None = None
+    actual_output_tokens: int | None = None
+    requested_output_tokens: int | None = None
+    finish_reason: str | None = None
+
+    @classmethod
+    def from_decision(
+        cls,
+        request: TraceRequest,
+        decision: RoutingDecision,
+        *,
+        mechanism_name: str,
+        execution_mode: str,
+        actual_kv_transfer_blocks: int,
+        effective_input_length: int | None,
+        cached_tokens: int,
+        session_reused: bool,
+        session_reset: bool,
+        latency_s: float | None,
+        ttft_s: float | None,
+        tpot_s: float | None,
+        prefill_request_priority: int | None = None,
+        decode_request_priority: int | None = None,
+        error: str | None = None,
+        actual_output_tokens: int | None = None,
+        requested_output_tokens: int | None = None,
+        finish_reason: str | None = None,
+    ) -> "RequestMetrics":
+        return cls(
+            request_id=request.request_id,
+            session_id=request.session_id,
+            turn_id=request.turn_id,
+            mechanism_name=mechanism_name,
+            execution_mode=execution_mode,
+            trace_timestamp_s=request.timestamp_s,
+            input_length=request.input_length,
+            output_length=request.output_length,
+            request_type=request.request_type,
+            policy_name=decision.policy_name,
+            assigned_prefill_node=decision.prefill_worker_id,
+            assigned_decode_node=decision.decode_worker_id,
+            assigned_decode_index=decision.decode_worker_index,
+            inflight_decode_load_at_assignment=decision.inflight_decode_load,
+            reuse_expected=decision.reuse_expected,
+            reuse_observed=decision.observed_reuse,
+            observed_overlap_blocks=decision.observed_overlap_blocks,
+            kv_transfer_blocks=decision.kv_transfer_blocks,
+            actual_kv_transfer_blocks=actual_kv_transfer_blocks,
+            cached_tokens=cached_tokens,
+            prefill_request_priority=prefill_request_priority,
+            decode_request_priority=decode_request_priority,
+            re_prefill_required=decision.re_prefill_required,
+            effective_input_length=effective_input_length,
+            session_reused=session_reused,
+            session_reset=session_reset,
+            latency_s=latency_s,
+            ttft_s=ttft_s,
+            tpot_s=tpot_s,
+            error=error,
+            actual_output_tokens=actual_output_tokens,
+            requested_output_tokens=requested_output_tokens,
+            finish_reason=finish_reason,
+        )
+
+
+def write_metrics_jsonl(path: Path, rows: list[RequestMetrics]) -> None:
+    path.parent.mkdir(parents=True, exist_ok=True)
+    with path.open("w", encoding="utf-8") as handle:
+        for row in rows:
+            handle.write(json.dumps(asdict(row), sort_keys=True) + "\n")
+
+
+def _is_failed_request(row: RequestMetrics) -> bool:
+    if row.error is not None:
+        return True
+    if row.finish_reason is not None:
+        fr = str(row.finish_reason).lower()
+        if "abort" in fr or "badrequest" in fr:
+            return True
+    return False
+
+
+def write_summary_json(
+    path: Path,
+    rows: list[RequestMetrics],
+    *,
+    trace_path: Path,
+    router_url: str | None,
+) -> None:
+    successful = [row for row in rows if not _is_failed_request(row)]
+    latencies = [row.latency_s for row in successful if row.latency_s is not None]
+    ttfts = [row.ttft_s for row in successful if row.ttft_s is not None]
+    tpots = [row.tpot_s for row in successful if row.tpot_s is not None]
+    per_decode_load = Counter(row.assigned_decode_node for row in rows)
+    per_prefill_load = Counter(row.assigned_prefill_node for row in rows)
+    prefill_priorities = Counter(
+        row.prefill_request_priority
+        for row in rows
+        if row.prefill_request_priority is not None
+    )
+    decode_priorities = Counter(
+        row.decode_request_priority
+        for row in rows
+        if row.decode_request_priority is not None
+    )
+
+    summary: dict[str, Any] = {
+        "trace_path": str(trace_path),
+        "router_url": router_url,
+        "request_count": len(rows),
+        "mechanisms": dict(sorted(Counter(row.mechanism_name for row in rows).items())),
+        "execution_modes": dict(sorted(Counter(row.execution_mode for row in rows).items())),
+        "latency_stats_s": _stats(latencies),
+        "ttft_stats_s": _stats(ttfts),
+        "tpot_stats_s": _stats(tpots),
+        "reuse_expected_count": sum(1 for row in rows if row.reuse_expected),
+        "reuse_observed_count": sum(1 for row in rows if row.reuse_observed),
+        "re_prefill_count": sum(1 for row in rows if row.re_prefill_required),
+        "cache_hit_request_count": sum(1 for row in rows if row.cached_tokens > 0),
+        "total_cached_tokens": sum(row.cached_tokens for row in rows),
+        "cached_tokens_stats": _stats([float(row.cached_tokens) for row in rows]),
+        "session_reused_count": sum(1 for row in rows if row.session_reused),
+        "session_reset_count": sum(1 for row in rows if row.session_reset),
+        "total_kv_transfer_blocks": sum(row.kv_transfer_blocks for row in rows),
+        "total_actual_kv_transfer_blocks": sum(
+            row.actual_kv_transfer_blocks for row in rows
+        ),
+        "per_decode_load": dict(sorted(per_decode_load.items())),
+        "per_prefill_load": dict(sorted(per_prefill_load.items())),
+        "prefill_request_priorities": {
+            str(key): value for key, value in sorted(prefill_priorities.items())
+        },
+        "decode_request_priorities": {
+            str(key): value for key, value in sorted(decode_priorities.items())
+        },
+        "error_count": sum(1 for row in rows if row.error is not None),
+        "abort_count": sum(
+            1
+            for row in rows
+            if row.error is None
+            and row.finish_reason is not None
+            and (
+                "abort" in str(row.finish_reason).lower()
+                or "badrequest" in str(row.finish_reason).lower()
+            )
+        ),
+        "failure_count": sum(1 for row in rows if _is_failed_request(row)),
+        "truncated_request_count": sum(
+            1
+            for row in rows
+            if row.actual_output_tokens is not None
+            and row.requested_output_tokens is not None
+            and row.requested_output_tokens > 1
+            and row.actual_output_tokens < row.requested_output_tokens * 0.5
+        ),
+        "actual_output_tokens_stats": _stats(
+            [float(row.actual_output_tokens) for row in rows if row.actual_output_tokens is not None]
+        ),
+    }
+    path.parent.mkdir(parents=True, exist_ok=True)
+    with path.open("w", encoding="utf-8") as handle:
+        json.dump(summary, handle, indent=2, sort_keys=True)
+
+
+def _stats(values: list[float | None]) -> dict[str, float] | None:
+    clean = [value for value in values if value is not None]
+    if not clean:
+        return None
+    clean.sort()
+    return {
+        "count": float(len(clean)),
+        "mean": statistics.fmean(clean),
+        "p50": _percentile(clean, 0.50),
+        "p90": _percentile(clean, 0.90),
+        "p99": _percentile(clean, 0.99),
+    }
+
+
+def _percentile(sorted_values: list[float], percentile: float) -> float:
+    if not sorted_values:
+        raise ValueError("sorted_values must not be empty")
+    if len(sorted_values) == 1:
+        return sorted_values[0]
+    index = round((len(sorted_values) - 1) * percentile)
+    return sorted_values[index]

src/agentic_pd_hybrid/microbench.py

@@ -0,0 +1,123 @@
+from __future__ import annotations
+
+import json
+from dataclasses import asdict, dataclass
+from math import ceil
+from pathlib import Path
+
+
+BLOCK_TOKEN_BUDGET = 24
+
+
+@dataclass(frozen=True)
+class SmallAppendTraceConfig:
+    output_path: Path
+    session_count: int = 8
+    turns_per_session: int = 3
+    initial_input_length: int = 10_000
+    append_input_length: int = 1_000
+    output_length: int = 1_000
+    inter_turn_gap_s: float = 1.0
+    session_stagger_s: float = 0.1
+    request_type: str = "coder"
+
+
+@dataclass(frozen=True)
+class SmallAppendTraceSummary:
+    output_path: str
+    session_count: int
+    turns_per_session: int
+    request_count: int
+    initial_input_length: int
+    append_input_length: int
+    output_length: int
+    inter_turn_gap_s: float
+    session_stagger_s: float
+
+
+def write_small_append_trace(config: SmallAppendTraceConfig) -> SmallAppendTraceSummary:
+    if config.session_count <= 0:
+        raise ValueError("session_count must be > 0")
+    if config.turns_per_session <= 0:
+        raise ValueError("turns_per_session must be > 0")
+    if config.initial_input_length < 0:
+        raise ValueError("initial_input_length must be >= 0")
+    if config.append_input_length < 0:
+        raise ValueError("append_input_length must be >= 0")
+    if config.output_length < 0:
+        raise ValueError("output_length must be >= 0")
+
+    config.output_path.parent.mkdir(parents=True, exist_ok=True)
+    records: list[dict[str, object]] = []
+    next_chat_id = 1_000_000
+
+    for session_idx in range(config.session_count):
+        root_chat_id = next_chat_id
+        previous_chat_id = -1
+        session_base_time = session_idx * config.session_stagger_s
+        base_block_count = ceil(config.initial_input_length / BLOCK_TOKEN_BUDGET)
+        base_hash_ids = [
+            _hash_id_for(session_idx=session_idx, block_idx=block_idx)
+            for block_idx in range(base_block_count)
+        ]
+
+        for turn_idx in range(config.turns_per_session):
+            chat_id = root_chat_id if turn_idx == 0 else next_chat_id
+            if turn_idx > 0:
+                next_chat_id += 1
+
+            input_length = config.initial_input_length + turn_idx * (
+                config.append_input_length + config.output_length
+            )
+            total_block_count = ceil(input_length / BLOCK_TOKEN_BUDGET)
+            hash_ids = base_hash_ids + [
+                _hash_id_for(
+                    session_idx=session_idx,
+                    block_idx=base_block_count + append_block_idx,
+                )
+                for append_block_idx in range(max(0, total_block_count - base_block_count))
+            ]
+
+            records.append(
+                {
+                    "chat_id": chat_id,
+                    "parent_chat_id": previous_chat_id,
+                    "timestamp": session_base_time
+                    + turn_idx * config.inter_turn_gap_s,
+                    "input_length": input_length,
+                    "output_length": config.output_length,
+                    "type": config.request_type,
+                    "turn": turn_idx + 1,
+                    "hash_ids": hash_ids,
+                }
+            )
+            previous_chat_id = chat_id
+
+        next_chat_id += 1
+
+    records.sort(key=lambda item: float(item["timestamp"]))
+    with config.output_path.open("w", encoding="utf-8") as handle:
+        for record in records:
+            handle.write(json.dumps(record, sort_keys=True) + "\n")
+
+    summary = SmallAppendTraceSummary(
+        output_path=str(config.output_path),
+        session_count=config.session_count,
+        turns_per_session=config.turns_per_session,
+        request_count=len(records),
+        initial_input_length=config.initial_input_length,
+        append_input_length=config.append_input_length,
+        output_length=config.output_length,
+        inter_turn_gap_s=config.inter_turn_gap_s,
+        session_stagger_s=config.session_stagger_s,
+    )
+    summary_path = config.output_path.with_suffix(
+        config.output_path.suffix + ".summary.json"
+    )
+    with summary_path.open("w", encoding="utf-8") as handle:
+        json.dump(asdict(summary), handle, indent=2, sort_keys=True)
+    return summary
+
+
+def _hash_id_for(*, session_idx: int, block_idx: int) -> int:
+    return session_idx * 1_000_000 + block_idx

src/agentic_pd_hybrid/pd_router.py

@@ -0,0 +1,466 @@
+from __future__ import annotations
+
+import argparse
+import asyncio
+import random
+import urllib.parse
+from dataclasses import dataclass
+from http import HTTPStatus
+from itertools import chain
+from typing import AsyncIterator
+
+import aiohttp
+import orjson
+import uvicorn
+from fastapi import FastAPI, HTTPException, Request
+from fastapi.responses import ORJSONResponse, Response, StreamingResponse
+
+_STREAM_CHUNK_SIZE = 1024 * 64
+
+
+@dataclass
+class RouterConfig:
+    host: str
+    port: int
+    prefill_urls: list[tuple[str, int]]
+    decode_urls: list[str]
+    prefill_policy: str = "round_robin"
+    decode_policy: str = "manual"
+    request_timeout_s: float = 1800.0
+
+
+class RouterState:
+    def __init__(self, config: RouterConfig):
+        if not config.prefill_urls:
+            raise ValueError("At least one prefill worker is required")
+        if not config.decode_urls:
+            raise ValueError("At least one decode worker is required")
+        self.config = config
+        self.prefill_cursor = 0
+        self.decode_cursor = 0
+        self.sticky_decode_map: dict[str, int] = {}
+
+    def select_pair(self, headers: dict[str, str]) -> tuple[str, int, str]:
+        prefill_url, bootstrap_port = self.config.prefill_urls[
+            self.prefill_cursor % len(self.config.prefill_urls)
+        ]
+        self.prefill_cursor += 1
+        decode_index = self._select_decode_index(headers)
+        return prefill_url, bootstrap_port, self.config.decode_urls[decode_index]
+
+    def _select_decode_index(self, headers: dict[str, str]) -> int:
+        target_worker = headers.get("x-smg-target-worker")
+        routing_key = headers.get("x-smg-routing-key")
+
+        if (
+            self.config.decode_policy == "consistent_hashing"
+            and target_worker is not None
+        ):
+            idx = int(target_worker)
+            if 0 <= idx < len(self.config.decode_urls):
+                return idx
+
+        if self.config.decode_policy == "manual" and routing_key:
+            cached = self.sticky_decode_map.get(routing_key)
+            if cached is not None:
+                return cached
+            idx = self.decode_cursor % len(self.config.decode_urls)
+            self.decode_cursor += 1
+            self.sticky_decode_map[routing_key] = idx
+            return idx
+
+        idx = self.decode_cursor % len(self.config.decode_urls)
+        self.decode_cursor += 1
+        return idx
+
+
+@dataclass
+class DpRouterConfig:
+    host: str
+    port: int
+    backend_urls: list[str]
+    backend_policy: str = "round_robin"
+    request_timeout_s: float = 1800.0
+
+
+class DpRouterState:
+    """DP (data-parallel) router: forward each request to exactly one backend."""
+
+    def __init__(self, config: DpRouterConfig):
+        if not config.backend_urls:
+            raise ValueError("At least one backend worker is required")
+        self.config = config
+        self.cursor = 0
+        self.sticky_map: dict[str, int] = {}
+
+    def select_backend(self, headers: dict[str, str]) -> str:
+        idx = self._select_index(headers)
+        return self.config.backend_urls[idx]
+
+    def _select_index(self, headers: dict[str, str]) -> int:
+        target_worker = headers.get("x-smg-target-worker")
+        routing_key = headers.get("x-smg-routing-key")
+
+        if (
+            self.config.backend_policy == "consistent_hashing"
+            and target_worker is not None
+        ):
+            idx = int(target_worker)
+            if 0 <= idx < len(self.config.backend_urls):
+                return idx
+
+        if self.config.backend_policy == "manual" and routing_key:
+            cached = self.sticky_map.get(routing_key)
+            if cached is not None:
+                return cached
+            idx = self.cursor % len(self.config.backend_urls)
+            self.cursor += 1
+            self.sticky_map[routing_key] = idx
+            return idx
+
+        idx = self.cursor % len(self.config.backend_urls)
+        self.cursor += 1
+        return idx
+
+
+app = FastAPI()
+router_state: RouterState | None = None
+dp_state: DpRouterState | None = None
+
+
+@app.get("/health")
+async def health() -> Response:
+    return Response(status_code=200)
+
+
+@app.get("/health_generate")
+async def health_generate() -> Response:
+    if dp_state is not None:
+        async with aiohttp.ClientSession() as session:
+            tasks = [
+                session.get(f"{url}/health_generate")
+                for url in dp_state.config.backend_urls
+            ]
+            for response in asyncio.as_completed(tasks):
+                async with await response:
+                    pass
+        return Response(status_code=200)
+    state = _require_state()
+    async with aiohttp.ClientSession() as session:
+        tasks = []
+        for server in chain(
+            (url for url, _ in state.config.prefill_urls),
+            state.config.decode_urls,
+        ):
+            tasks.append(session.get(f"{server}/health_generate"))
+        for response in asyncio.as_completed(tasks):
+            async with await response:
+                pass
+    return Response(status_code=200)
+
+
+@app.get("/v1/models")
+async def models() -> ORJSONResponse:
+    if dp_state is not None:
+        async with aiohttp.ClientSession() as session:
+            async with session.get(f"{dp_state.config.backend_urls[0]}/v1/models") as resp:
+                payload = await resp.json()
+                return ORJSONResponse(payload, status_code=resp.status)
+    state = _require_state()
+    async with aiohttp.ClientSession() as session:
+        async with session.get(f"{state.config.prefill_urls[0][0]}/v1/models") as response:
+            payload = await response.json()
+            return ORJSONResponse(payload, status_code=response.status)
+
+
+@app.post("/v1/chat/completions")
+async def chat_completions(request: Request) -> Response:
+    request_data = await request.json()
+    headers = {key.lower(): value for key, value in request.headers.items()}
+    return await _forward_to_backend(
+        request_data=request_data,
+        headers=headers,
+        endpoint_name="v1/chat/completions",
+    )
+
+
+@app.post("/v1/completions")
+async def completions(request: Request) -> Response:
+    request_data = await request.json()
+    headers = {key.lower(): value for key, value in request.headers.items()}
+    return await _forward_to_backend(
+        request_data=request_data,
+        headers=headers,
+        endpoint_name="v1/completions",
+    )
+
+
+@app.post("/generate")
+async def generate(request: Request) -> Response:
+    request_data = await request.json()
+    headers = {key.lower(): value for key, value in request.headers.items()}
+    return await _forward_to_backend(
+        request_data=request_data,
+        headers=headers,
+        endpoint_name="generate",
+    )
+
+
+async def _forward_to_backend(
+    *,
+    request_data: dict,
+    headers: dict[str, str],
+    endpoint_name: str,
+) -> Response:
+    # DP mode: forward to a single backend
+    if dp_state is not None:
+        return await _forward_to_dp_backend(
+            request_data=request_data,
+            headers=headers,
+            endpoint_name=endpoint_name,
+        )
+
+    # PD mode: coordinate prefill + decode
+    state = _require_state()
+    prefill_server, bootstrap_port, decode_server = state.select_pair(headers)
+    prefill_request, decode_request = _build_backend_requests(
+        request_data=request_data,
+        prefill_server=prefill_server,
+        bootstrap_port=bootstrap_port,
+    )
+
+    if request_data.get("stream", False):
+        return StreamingResponse(
+            _stream_generate(
+                prefill_request=prefill_request,
+                decode_request=decode_request,
+                prefill_server=prefill_server,
+                decode_server=decode_server,
+                endpoint_name=endpoint_name,
+                timeout_s=state.config.request_timeout_s,
+            ),
+            media_type="text/event-stream",
+        )
+
+    async with aiohttp.ClientSession(
+        timeout=aiohttp.ClientTimeout(total=state.config.request_timeout_s)
+    ) as session:
+        prefill_response, decode_response = await asyncio.gather(
+            session.post(f"{prefill_server}/{endpoint_name}", json=prefill_request),
+            session.post(f"{decode_server}/{endpoint_name}", json=decode_request),
+        )
+        async with prefill_response:
+            await prefill_response.read()
+        async with decode_response:
+            body = await decode_response.read()
+            return Response(
+                content=body,
+                status_code=decode_response.status,
+                media_type=decode_response.content_type,
+            )
+
+
+async def _forward_to_dp_backend(
+    *,
+    request_data: dict,
+    headers: dict[str, str],
+    endpoint_name: str,
+) -> Response:
+    assert dp_state is not None
+    backend_server = dp_state.select_backend(headers)
+    cleaned = _strip_internal_fields(request_data)
+    timeout_s = dp_state.config.request_timeout_s
+
+    if request_data.get("stream", False):
+        return StreamingResponse(
+            _stream_dp_generate(
+                request_data=cleaned,
+                backend_server=backend_server,
+                endpoint_name=endpoint_name,
+                timeout_s=timeout_s,
+            ),
+            media_type="text/event-stream",
+        )
+
+    async with aiohttp.ClientSession(
+        timeout=aiohttp.ClientTimeout(total=timeout_s)
+    ) as session:
+        async with session.post(
+            f"{backend_server}/{endpoint_name}", json=cleaned
+        ) as response:
+            body = await response.read()
+            return Response(
+                content=body,
+                status_code=response.status,
+                media_type=response.content_type,
+            )
+
+
+async def _stream_dp_generate(
+    *,
+    request_data: dict,
+    backend_server: str,
+    endpoint_name: str,
+    timeout_s: float,
+) -> AsyncIterator[bytes]:
+    async with aiohttp.ClientSession(
+        timeout=aiohttp.ClientTimeout(total=timeout_s)
+    ) as session:
+        async with session.post(
+            f"{backend_server}/{endpoint_name}", json=request_data
+        ) as response:
+            if response.status != HTTPStatus.OK:
+                payload = await response.read()
+                yield payload
+                return
+            async for chunk in response.content.iter_chunked(_STREAM_CHUNK_SIZE):
+                yield chunk
+
+
+async def _stream_generate(
+    *,
+    prefill_request: dict,
+    decode_request: dict,
+    prefill_server: str,
+    decode_server: str,
+    endpoint_name: str,
+    timeout_s: float,
+) -> AsyncIterator[bytes]:
+    async with aiohttp.ClientSession(
+        timeout=aiohttp.ClientTimeout(total=timeout_s)
+    ) as session:
+        prefill_response, decode_response = await asyncio.gather(
+            session.post(f"{prefill_server}/{endpoint_name}", json=prefill_request),
+            session.post(f"{decode_server}/{endpoint_name}", json=decode_request),
+        )
+        async with prefill_response, decode_response:
+            if decode_response.status != HTTPStatus.OK:
+                payload = await decode_response.read()
+                yield payload
+                return
+            async for chunk in decode_response.content.iter_chunked(_STREAM_CHUNK_SIZE):
+                yield chunk
+
+
+def _build_bootstrap_payload(prefill_server: str, bootstrap_port: int) -> dict[str, object]:
+    parsed_url = urllib.parse.urlparse(prefill_server)
+    hostname = parsed_url.hostname
+    if hostname is None:
+        raise HTTPException(
+            status_code=500,
+            detail=f"Unable to parse prefill hostname from {prefill_server}",
+        )
+    return {
+        "bootstrap_host": hostname,
+        "bootstrap_port": bootstrap_port,
+        "bootstrap_room": random.randint(0, 2**63 - 1),
+    }
+
+
+def _build_backend_requests(
+    *,
+    request_data: dict,
+    prefill_server: str,
+    bootstrap_port: int,
+) -> tuple[dict, dict]:
+    prefill_priority = request_data.get("smg_prefill_priority")
+    decode_priority = request_data.get("smg_decode_priority")
+
+    prefill_request = _strip_internal_fields(request_data)
+    decode_request = _strip_internal_fields(request_data)
+    bootstrap_payload = _build_bootstrap_payload(prefill_server, bootstrap_port)
+    prefill_request.update(bootstrap_payload)
+    decode_request.update(bootstrap_payload)
+
+    # session_params is only meaningful for the decode worker (streaming session
+    # KV reuse).  Sending it to the prefill worker causes the D side to
+    # short-circuit with local-prefill on already-open sessions, returning
+    # truncated responses while P's KV transfer gets aborted.
+    prefill_request.pop("session_params", None)
+
+    if prefill_priority is not None:
+        prefill_request["priority"] = int(prefill_priority)
+    if decode_priority is not None:
+        decode_request["priority"] = int(decode_priority)
+    return prefill_request, decode_request
+
+
+def _strip_internal_fields(request_data: dict) -> dict:
+    cleaned = request_data.copy()
+    cleaned.pop("smg_prefill_priority", None)
+    cleaned.pop("smg_decode_priority", None)
+    return cleaned
+
+
+def _require_state() -> RouterState:
+    if router_state is None:
+        raise HTTPException(status_code=500, detail="router not initialized")
+    return router_state
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser(description="Minimal local PD / DP router")
+    parser.add_argument("--host", default="127.0.0.1")
+    parser.add_argument("--port", type=int, default=8000)
+    parser.add_argument(
+        "--prefill",
+        nargs=2,
+        metavar=("URL", "BOOTSTRAP_PORT"),
+        action="append",
+        default=None,
+    )
+    parser.add_argument(
+        "--decode",
+        action="append",
+        default=None,
+    )
+    parser.add_argument("--prefill-policy", default="round_robin")
+    parser.add_argument("--decode-policy", default="manual")
+    parser.add_argument(
+        "--backend",
+        action="append",
+        default=None,
+        help="Backend URL for DP (data-parallel) mode. Repeat for each worker.",
+    )
+    parser.add_argument(
+        "--backend-policy",
+        default="round_robin",
+        help="Routing policy for DP mode: round_robin, manual, consistent_hashing.",
+    )
+    parser.add_argument("--request-timeout-s", type=float, default=1800.0)
+    args = parser.parse_args()
+
+    global router_state, dp_state
+
+    if args.backend:
+        # DP mode: simple forward to one of N backends
+        dp_state = DpRouterState(
+            DpRouterConfig(
+                host=args.host,
+                port=args.port,
+                backend_urls=list(args.backend),
+                backend_policy=args.backend_policy,
+                request_timeout_s=args.request_timeout_s,
+            )
+        )
+    elif args.prefill and args.decode:
+        # PD mode: prefill/decode coordination
+        router_state = RouterState(
+            RouterConfig(
+                host=args.host,
+                port=args.port,
+                prefill_urls=[(url, int(port)) for url, port in args.prefill],
+                decode_urls=list(args.decode),
+                prefill_policy=args.prefill_policy,
+                decode_policy=args.decode_policy,
+                request_timeout_s=args.request_timeout_s,
+            )
+        )
+    else:
+        parser.error("Either --backend (DP mode) or both --prefill and --decode (PD mode) are required")
+
+    uvicorn.run(app, host=args.host, port=args.port, log_level="info")
+
+
+if __name__ == "__main__":
+    main()

src/agentic_pd_hybrid/policies.py

@@ -0,0 +1,271 @@
+from __future__ import annotations
+
+from collections import Counter
+from dataclasses import dataclass, field
+from typing import Protocol
+
+from agentic_pd_hybrid.topology import SingleNodeTopology
+from agentic_pd_hybrid.trace import TraceRequest
+
+
+@dataclass
+class SessionRouteState:
+    last_decode_worker: str | None = None
+
+
+@dataclass
+class RoutingDecision:
+    policy_name: str
+    prefill_worker_id: str
+    decode_worker_id: str
+    decode_worker_index: int
+    reuse_expected: bool
+    observed_overlap_blocks: int
+    kv_transfer_blocks: int
+    inflight_decode_load: int
+    session_id: str
+    request_id: str
+    turn_id: int
+
+    @property
+    def observed_reuse(self) -> bool:
+        return self.observed_overlap_blocks > 0
+
+    @property
+    def re_prefill_required(self) -> bool:
+        return self.turn_id > 1 and self.observed_overlap_blocks == 0
+
+
+@dataclass
+class RoutingState:
+    prefill_cursor: int = 0
+    decode_cursor: int = 0
+    session_state: dict[str, SessionRouteState] = field(default_factory=dict)
+    inflight_decode: Counter[str] = field(default_factory=Counter)
+    decode_assignment_counts: Counter[str] = field(default_factory=Counter)
+    decode_resident_blocks: dict[str, set[int]] = field(default_factory=dict)
+    # Migration support: per-(session_id, decode_worker_id) admission reject counter.
+    # KvAwarePolicy uses this to skip D's that have repeatedly rejected this session
+    # (avoids the structural starvation observed in TEAM_REPORT §2.1).
+    session_d_rejects: Counter[tuple[str, str]] = field(default_factory=Counter)
+
+    @classmethod
+    def create(cls, topology: SingleNodeTopology) -> "RoutingState":
+        return cls(
+            decode_resident_blocks={
+                worker.worker_id: set() for worker in topology.route_workers
+            }
+        )
+
+    def next_prefill_worker_id(self, topology: SingleNodeTopology) -> str:
+        if not topology.prefill_workers:
+            return "none"
+        worker = topology.prefill_workers[self.prefill_cursor % len(topology.prefill_workers)]
+        self.prefill_cursor += 1
+        return worker.worker_id
+
+    def next_decode_worker_id(self, topology: SingleNodeTopology) -> str:
+        route_workers = topology.route_workers
+        worker = route_workers[self.decode_cursor % len(route_workers)]
+        self.decode_cursor += 1
+        return worker.worker_id
+
+    def record_admission_reject(self, session_id: str, decode_worker_id: str) -> int:
+        """Increment per-(session, D) rejection counter. Returns new count."""
+        key = (session_id, decode_worker_id)
+        self.session_d_rejects[key] += 1
+        return self.session_d_rejects[key]
+
+    def finish(self, request: TraceRequest, decision: RoutingDecision) -> None:
+        session = self.session_state.setdefault(request.session_id, SessionRouteState())
+        session.last_decode_worker = decision.decode_worker_id
+        self.decode_resident_blocks[decision.decode_worker_id].update(request.hash_ids)
+        self.inflight_decode[decision.decode_worker_id] -= 1
+        if self.inflight_decode[decision.decode_worker_id] <= 0:
+            del self.inflight_decode[decision.decode_worker_id]
+
+
+class RoutingPolicy(Protocol):
+    name: str
+
+    def select(
+        self,
+        request: TraceRequest,
+        *,
+        topology: SingleNodeTopology,
+        state: RoutingState,
+    ) -> RoutingDecision:
+        ...
+
+
+@dataclass(frozen=True)
+class DefaultPolicy:
+    name: str = "default"
+
+    def select(
+        self,
+        request: TraceRequest,
+        *,
+        topology: SingleNodeTopology,
+        state: RoutingState,
+    ) -> RoutingDecision:
+        prefill_worker_id = state.next_prefill_worker_id(topology)
+        decode_worker_id = state.next_decode_worker_id(topology)
+        return _build_decision(
+            policy_name=self.name,
+            request=request,
+            topology=topology,
+            state=state,
+            prefill_worker_id=prefill_worker_id,
+            decode_worker_id=decode_worker_id,
+            reuse_expected=False,
+        )
+
+
+@dataclass(frozen=True)
+class StickyDecodePolicy:
+    name: str = "sticky"
+
+    def select(
+        self,
+        request: TraceRequest,
+        *,
+        topology: SingleNodeTopology,
+        state: RoutingState,
+    ) -> RoutingDecision:
+        session = state.session_state.get(request.session_id)
+        prefill_worker_id = state.next_prefill_worker_id(topology)
+        if request.turn_id > 1 and session and session.last_decode_worker is not None:
+            decode_worker_id = session.last_decode_worker
+            reuse_expected = True
+        else:
+            decode_worker_id = state.next_decode_worker_id(topology)
+            reuse_expected = False
+        return _build_decision(
+            policy_name=self.name,
+            request=request,
+            topology=topology,
+            state=state,
+            prefill_worker_id=prefill_worker_id,
+            decode_worker_id=decode_worker_id,
+            reuse_expected=reuse_expected,
+        )
+
+
+@dataclass(frozen=True)
+class KvAwarePolicy:
+    name: str = "kv-aware"
+    sticky_bonus: int = 1
+    # Session migration: when (session, D) has been rejected this many times,
+    # skip D entirely for this session (force migration to another D).
+    # 0 disables the mechanism. Default 3 picked empirically to allow brief
+    # transient saturation without panicking, but to reroute persistent starvation.
+    migration_reject_threshold: int = 3
+
+    def select(
+        self,
+        request: TraceRequest,
+        *,
+        topology: SingleNodeTopology,
+        state: RoutingState,
+    ) -> RoutingDecision:
+        prefill_worker_id = state.next_prefill_worker_id(topology)
+        session = state.session_state.get(request.session_id)
+
+        best_decode_worker_id: str | None = None
+        best_score: tuple[int, int, int, int] | None = None
+        candidates_considered = 0
+        for worker in topology.route_workers:
+            # Migration: skip workers that have rejected this session too many times.
+            # If all candidates get filtered (degenerate case), fall through to
+            # un-filtered selection below.
+            if self.migration_reject_threshold > 0:
+                rejects = state.session_d_rejects.get(
+                    (request.session_id, worker.worker_id), 0
+                )
+                if rejects >= self.migration_reject_threshold:
+                    continue
+            candidates_considered += 1
+            overlap = _overlap_blocks(request, state, worker.worker_id)
+            sticky = int(session is not None and session.last_decode_worker == worker.worker_id)
+            inflight_penalty = -state.inflight_decode.get(worker.worker_id, 0)
+            assignment_penalty = -state.decode_assignment_counts.get(worker.worker_id, 0)
+            score = (
+                overlap + sticky * self.sticky_bonus,
+                sticky,
+                inflight_penalty,
+                assignment_penalty,
+            )
+            if best_score is None or score > best_score:
+                best_score = score
+                best_decode_worker_id = worker.worker_id
+
+        # Degenerate fallback: every D was filtered. Pick the least-rejected D.
+        if best_decode_worker_id is None:
+            best_decode_worker_id = min(
+                (w.worker_id for w in topology.route_workers),
+                key=lambda wid: state.session_d_rejects.get(
+                    (request.session_id, wid), 0
+                ),
+            )
+            best_score = (0, 0, 0, 0)
+
+        assert best_decode_worker_id is not None
+        reuse_expected = bool(best_score and best_score[0] > 0)
+        return _build_decision(
+            policy_name=self.name,
+            request=request,
+            topology=topology,
+            state=state,
+            prefill_worker_id=prefill_worker_id,
+            decode_worker_id=best_decode_worker_id,
+            reuse_expected=reuse_expected,
+        )
+
+
+def create_policy(name: str, *, migration_reject_threshold: int = 3) -> RoutingPolicy:
+    normalized = name.strip().lower()
+    if normalized == "default":
+        return DefaultPolicy()
+    if normalized == "sticky":
+        return StickyDecodePolicy()
+    if normalized in {"kv-aware", "kv_aware", "kv"}:
+        return KvAwarePolicy(migration_reject_threshold=migration_reject_threshold)
+    raise ValueError(f"Unsupported policy: {name}")
+
+
+def _build_decision(
+    *,
+    policy_name: str,
+    request: TraceRequest,
+    topology: SingleNodeTopology,
+    state: RoutingState,
+    prefill_worker_id: str,
+    decode_worker_id: str,
+    reuse_expected: bool,
+) -> RoutingDecision:
+    overlap = _overlap_blocks(request, state, decode_worker_id)
+    state.inflight_decode[decode_worker_id] += 1
+    state.decode_assignment_counts[decode_worker_id] += 1
+    return RoutingDecision(
+        policy_name=policy_name,
+        prefill_worker_id=prefill_worker_id,
+        decode_worker_id=decode_worker_id,
+        decode_worker_index=topology.route_index(decode_worker_id),
+        reuse_expected=reuse_expected,
+        observed_overlap_blocks=overlap,
+        kv_transfer_blocks=max(0, len(request.hash_ids) - overlap),
+        inflight_decode_load=state.inflight_decode[decode_worker_id],
+        session_id=request.session_id,
+        request_id=request.request_id,
+        turn_id=request.turn_id,
+    )
+
+
+def _overlap_blocks(
+    request: TraceRequest,
+    state: RoutingState,
+    decode_worker_id: str,
+) -> int:
+    resident = state.decode_resident_blocks.get(decode_worker_id, set())
+    return sum(1 for block in request.hash_ids if block in resident)

src/agentic_pd_hybrid/profile.py

@@ -0,0 +1,511 @@
+from __future__ import annotations
+
+import json
+import statistics
+from collections import Counter, defaultdict
+from dataclasses import dataclass
+from pathlib import Path
+from typing import Any, Iterable
+
+from agentic_pd_hybrid.trace import TraceRequest, load_trace
+
+
+@dataclass(frozen=True)
+class ProfileConfig:
+    trace_path: Path
+    output_path: Path | None = None
+    metrics_path: Path | None = None
+    baseline_metrics_path: Path | None = None
+    candidate_metrics_path: Path | None = None
+    direct_max_uncached_tokens: int = 2048
+
+
+def write_profile(config: ProfileConfig) -> dict[str, Any]:
+    report = build_profile(config)
+    if config.output_path is not None:
+        config.output_path.parent.mkdir(parents=True, exist_ok=True)
+        with config.output_path.open("w", encoding="utf-8") as handle:
+            json.dump(report, handle, indent=2, sort_keys=True)
+    return report
+
+
+def build_profile(config: ProfileConfig) -> dict[str, Any]:
+    requests = load_trace(config.trace_path)
+    features = _build_trace_features(
+        requests,
+        direct_max_uncached_tokens=config.direct_max_uncached_tokens,
+    )
+    report: dict[str, Any] = {
+        "trace_path": str(config.trace_path),
+        "direct_max_uncached_tokens": config.direct_max_uncached_tokens,
+        "trace_profile": _trace_profile(requests, features),
+    }
+
+    if config.metrics_path is not None:
+        metrics = _load_jsonl(config.metrics_path)
+        report["metrics_path"] = str(config.metrics_path)
+        report["metrics_profile"] = _metrics_profile(metrics, features)
+
+    if (
+        config.baseline_metrics_path is not None
+        and config.candidate_metrics_path is not None
+    ):
+        baseline = _load_jsonl(config.baseline_metrics_path)
+        candidate = _load_jsonl(config.candidate_metrics_path)
+        report["baseline_metrics_path"] = str(config.baseline_metrics_path)
+        report["candidate_metrics_path"] = str(config.candidate_metrics_path)
+        report["baseline_profile"] = _metrics_profile(baseline, features)
+        report["candidate_profile"] = _metrics_profile(candidate, features)
+        report["paired_comparison"] = _paired_comparison(
+            baseline=baseline,
+            candidate=candidate,
+            features=features,
+        )
+
+    return report
+
+
+def print_profile_summary(report: dict[str, Any]) -> None:
+    trace = report["trace_profile"]
+    print(
+        "trace: "
+        f"{trace['request_count']} requests, "
+        f"{trace['session_count']} sessions, "
+        f"{trace['multi_turn_session_count']} multi-turn sessions"
+    )
+    print(
+        "direct-eligible turns: "
+        f"{trace['direct_eligible_turn2plus_count']}/"
+        f"{trace['turn2plus_count']} "
+        f"({trace['direct_eligible_turn2plus_ratio']:.3f})"
+    )
+    append_stats = trace.get("append_input_tokens_stats")
+    output_stats = trace.get("output_tokens_stats")
+    if append_stats is not None:
+        print(
+            "append tokens: "
+            f"mean={append_stats['mean']:.1f} "
+            f"p50={append_stats['p50']:.1f} "
+            f"p90={append_stats['p90']:.1f} "
+            f"p99={append_stats['p99']:.1f}"
+        )
+    if output_stats is not None:
+        print(
+            "output tokens: "
+            f"mean={output_stats['mean']:.1f} "
+            f"p50={output_stats['p50']:.1f} "
+            f"p90={output_stats['p90']:.1f} "
+            f"p99={output_stats['p99']:.1f}"
+        )
+
+    comparison = report.get("paired_comparison")
+    if isinstance(comparison, dict):
+        overall = comparison.get("overall", {})
+        delta = overall.get("latency_delta_s_stats")
+        if delta is not None:
+            print(
+                "candidate - baseline E2E: "
+                f"mean={delta['mean']:.3f}s "
+                f"p50={delta['p50']:.3f}s "
+                f"p90={delta['p90']:.3f}s"
+            )
+        print(
+            "paired wins/losses: "
+            f"{overall.get('candidate_faster_count', 0)} faster, "
+            f"{overall.get('candidate_slower_count', 0)} slower, "
+            f"{overall.get('paired_count', 0)} paired"
+        )
+
+
+@dataclass(frozen=True)
+class _TraceFeature:
+    request_id: str
+    session_id: str
+    turn_id: int
+    input_length: int
+    output_length: int
+    resident_tokens: int
+    append_input_tokens: int | None
+    inter_turn_gap_s: float | None
+    overlap_blocks_with_previous: int | None
+    overlap_ratio_with_previous: float | None
+    direct_eligible: bool
+    turn_class: str
+    append_bin: str
+    input_bin: str
+    output_bin: str
+    resident_bin: str
+
+
+def _build_trace_features(
+    requests: list[TraceRequest],
+    *,
+    direct_max_uncached_tokens: int,
+) -> dict[str, _TraceFeature]:
+    ordered_by_session: dict[str, list[TraceRequest]] = defaultdict(list)
+    for request in requests:
+        ordered_by_session[request.session_id].append(request)
+
+    previous_by_request_id: dict[str, TraceRequest | None] = {}
+    for session_requests in ordered_by_session.values():
+        ordered = sorted(
+            session_requests,
+            key=lambda request: (request.timestamp_s, request.turn_id, request.chat_id),
+        )
+        previous: TraceRequest | None = None
+        for request in ordered:
+            previous_by_request_id[request.request_id] = previous
+            previous = request
+
+    features: dict[str, _TraceFeature] = {}
+    for request in requests:
+        previous = previous_by_request_id.get(request.request_id)
+        append_input_tokens: int | None = None
+        inter_turn_gap_s: float | None = None
+        overlap_blocks: int | None = None
+        overlap_ratio: float | None = None
+        direct_eligible = False
+
+        if previous is not None:
+            append_input_tokens = request.input_length - (
+                previous.input_length + previous.output_length
+            )
+            inter_turn_gap_s = request.timestamp_s - previous.timestamp_s
+            previous_blocks = set(previous.hash_ids)
+            overlap_blocks = sum(1 for block in request.hash_ids if block in previous_blocks)
+            overlap_ratio = (
+                overlap_blocks / len(request.hash_ids) if request.hash_ids else 0.0
+            )
+            direct_eligible = (
+                append_input_tokens > 0
+                and append_input_tokens <= direct_max_uncached_tokens
+                and overlap_blocks > 0
+            )
+
+        features[request.request_id] = _TraceFeature(
+            request_id=request.request_id,
+            session_id=request.session_id,
+            turn_id=request.turn_id,
+            input_length=request.input_length,
+            output_length=request.output_length,
+            resident_tokens=request.input_length + request.output_length,
+            append_input_tokens=append_input_tokens,
+            inter_turn_gap_s=inter_turn_gap_s,
+            overlap_blocks_with_previous=overlap_blocks,
+            overlap_ratio_with_previous=overlap_ratio,
+            direct_eligible=direct_eligible,
+            turn_class="turn1" if request.turn_id <= 1 else "turn2plus",
+            append_bin=_token_bin(append_input_tokens),
+            input_bin=_token_bin(request.input_length),
+            output_bin=_token_bin(request.output_length),
+            resident_bin=_token_bin(request.input_length + request.output_length),
+        )
+    return features
+
+
+def _trace_profile(
+    requests: list[TraceRequest],
+    features: dict[str, _TraceFeature],
+) -> dict[str, Any]:
+    session_turns = Counter(request.session_id for request in requests)
+    turn2plus = [feature for feature in features.values() if feature.turn_id > 1]
+    direct_eligible = [feature for feature in turn2plus if feature.direct_eligible]
+    append_values = [
+        feature.append_input_tokens
+        for feature in turn2plus
+        if feature.append_input_tokens is not None
+    ]
+    positive_append_values = [
+        value for value in append_values if value is not None and value > 0
+    ]
+    overlap_ratios = [
+        feature.overlap_ratio_with_previous
+        for feature in turn2plus
+        if feature.overlap_ratio_with_previous is not None
+    ]
+    gaps = [
+        feature.inter_turn_gap_s
+        for feature in turn2plus
+        if feature.inter_turn_gap_s is not None
+    ]
+    return {
+        "request_count": len(requests),
+        "session_count": len(session_turns),
+        "multi_turn_session_count": sum(1 for turns in session_turns.values() if turns > 1),
+        "turn2plus_count": len(turn2plus),
+        "direct_eligible_turn2plus_count": len(direct_eligible),
+        "direct_eligible_turn2plus_ratio": (
+            len(direct_eligible) / len(turn2plus) if turn2plus else 0.0
+        ),
+        "turn_count_distribution": dict(sorted(Counter(session_turns.values()).items())),
+        "request_type_distribution": dict(
+            sorted(Counter(request.request_type for request in requests).items())
+        ),
+        "turn_id_distribution": dict(
+            sorted(Counter(request.turn_id for request in requests).items())
+        ),
+        "append_bin_distribution": dict(
+            sorted(Counter(feature.append_bin for feature in turn2plus).items())
+        ),
+        "input_bin_distribution": dict(
+            sorted(Counter(feature.input_bin for feature in features.values()).items())
+        ),
+        "output_bin_distribution": dict(
+            sorted(Counter(feature.output_bin for feature in features.values()).items())
+        ),
+        "resident_bin_distribution": dict(
+            sorted(Counter(feature.resident_bin for feature in features.values()).items())
+        ),
+        "input_tokens_stats": _stats(
+            [float(request.input_length) for request in requests]
+        ),
+        "output_tokens_stats": _stats(
+            [float(request.output_length) for request in requests]
+        ),
+        "resident_tokens_stats": _stats(
+            [float(feature.resident_tokens) for feature in features.values()]
+        ),
+        "append_input_tokens_stats": _stats(
+            [float(value) for value in append_values if value is not None]
+        ),
+        "positive_append_input_tokens_stats": _stats(
+            [float(value) for value in positive_append_values]
+        ),
+        "inter_turn_gap_s_stats": _stats([float(value) for value in gaps]),
+        "overlap_ratio_stats": _stats([float(value) for value in overlap_ratios]),
+        "non_positive_append_count": sum(
+            1 for value in append_values if value is not None and value <= 0
+        ),
+    }
+
+
+def _metrics_profile(
+    rows: list[dict[str, Any]],
+    features: dict[str, _TraceFeature],
+) -> dict[str, Any]:
+    return {
+        "request_count": len(rows),
+        "mechanism_distribution": dict(
+            sorted(Counter(str(row.get("mechanism_name")) for row in rows).items())
+        ),
+        "execution_mode_distribution": dict(
+            sorted(Counter(str(row.get("execution_mode")) for row in rows).items())
+        ),
+        "latency_s_stats": _stats(_numeric_values(rows, "latency_s")),
+        "ttft_s_stats": _stats(_numeric_values(rows, "ttft_s")),
+        "tpot_s_stats": _stats(_numeric_values(rows, "tpot_s")),
+        "cached_tokens_stats": _stats(_numeric_values(rows, "cached_tokens")),
+        "actual_kv_transfer_blocks_stats": _stats(
+            _numeric_values(rows, "actual_kv_transfer_blocks")
+        ),
+        "session_reused_count": sum(1 for row in rows if row.get("session_reused")),
+        "session_reset_count": sum(1 for row in rows if row.get("session_reset")),
+        "error_count": sum(1 for row in rows if row.get("error") is not None),
+        "by_turn_class": _group_metrics(rows, features, lambda feature, _row: feature.turn_class),
+        "by_direct_eligible": _group_metrics(
+            rows,
+            features,
+            lambda feature, _row: "eligible" if feature.direct_eligible else "not_eligible",
+        ),
+        "by_append_bin": _group_metrics(rows, features, lambda feature, _row: feature.append_bin),
+        "by_resident_bin": _group_metrics(
+            rows,
+            features,
+            lambda feature, _row: feature.resident_bin,
+        ),
+        "by_execution_mode": _group_metrics(
+            rows,
+            features,
+            lambda _feature, row: str(row.get("execution_mode")),
+        ),
+    }
+
+
+def _paired_comparison(
+    *,
+    baseline: list[dict[str, Any]],
+    candidate: list[dict[str, Any]],
+    features: dict[str, _TraceFeature],
+) -> dict[str, Any]:
+    baseline_by_id = {
+        str(row.get("request_id")): row
+        for row in baseline
+        if row.get("latency_s") is not None
+    }
+    candidate_by_id = {
+        str(row.get("request_id")): row
+        for row in candidate
+        if row.get("latency_s") is not None
+    }
+    paired_ids = sorted(set(baseline_by_id) & set(candidate_by_id))
+    pairs = [
+        (baseline_by_id[request_id], candidate_by_id[request_id], features.get(request_id))
+        for request_id in paired_ids
+    ]
+    pairs = [pair for pair in pairs if pair[2] is not None]
+
+    return {
+        "overall": _delta_summary(pairs),
+        "by_turn_class": _group_deltas(
+            pairs,
+            lambda feature, _base, _cand: feature.turn_class,
+        ),
+        "by_direct_eligible": _group_deltas(
+            pairs,
+            lambda feature, _base, _cand: (
+                "eligible" if feature.direct_eligible else "not_eligible"
+            ),
+        ),
+        "by_append_bin": _group_deltas(
+            pairs,
+            lambda feature, _base, _cand: feature.append_bin,
+        ),
+        "by_resident_bin": _group_deltas(
+            pairs,
+            lambda feature, _base, _cand: feature.resident_bin,
+        ),
+        "by_candidate_execution_mode": _group_deltas(
+            pairs,
+            lambda _feature, _base, cand: str(cand.get("execution_mode")),
+        ),
+    }
+
+
+def _group_metrics(
+    rows: list[dict[str, Any]],
+    features: dict[str, _TraceFeature],
+    key_fn,
+) -> dict[str, Any]:
+    grouped: dict[str, list[dict[str, Any]]] = defaultdict(list)
+    for row in rows:
+        feature = features.get(str(row.get("request_id")))
+        if feature is None:
+            continue
+        grouped[str(key_fn(feature, row))].append(row)
+    return {
+        key: {
+            "count": len(group_rows),
+            "latency_s_stats": _stats(_numeric_values(group_rows, "latency_s")),
+            "ttft_s_stats": _stats(_numeric_values(group_rows, "ttft_s")),
+            "tpot_s_stats": _stats(_numeric_values(group_rows, "tpot_s")),
+            "session_reused_count": sum(
+                1 for row in group_rows if row.get("session_reused")
+            ),
+            "error_count": sum(1 for row in group_rows if row.get("error") is not None),
+        }
+        for key, group_rows in sorted(grouped.items())
+    }
+
+
+def _group_deltas(
+    pairs: list[tuple[dict[str, Any], dict[str, Any], _TraceFeature | None]],
+    key_fn,
+) -> dict[str, Any]:
+    grouped: dict[str, list[tuple[dict[str, Any], dict[str, Any], _TraceFeature | None]]] = (
+        defaultdict(list)
+    )
+    for base, cand, feature in pairs:
+        if feature is None:
+            continue
+        grouped[str(key_fn(feature, base, cand))].append((base, cand, feature))
+    return {key: _delta_summary(group_pairs) for key, group_pairs in sorted(grouped.items())}
+
+
+def _delta_summary(
+    pairs: list[tuple[dict[str, Any], dict[str, Any], _TraceFeature | None]],
+) -> dict[str, Any]:
+    latency_deltas = [
+        float(cand["latency_s"]) - float(base["latency_s"])
+        for base, cand, _feature in pairs
+        if base.get("latency_s") is not None and cand.get("latency_s") is not None
+    ]
+    ttft_deltas = [
+        float(cand["ttft_s"]) - float(base["ttft_s"])
+        for base, cand, _feature in pairs
+        if base.get("ttft_s") is not None and cand.get("ttft_s") is not None
+    ]
+    return {
+        "paired_count": len(latency_deltas),
+        "candidate_faster_count": sum(1 for delta in latency_deltas if delta < 0),
+        "candidate_slower_count": sum(1 for delta in latency_deltas if delta > 0),
+        "latency_delta_s_stats": _stats(latency_deltas),
+        "ttft_delta_s_stats": _stats(ttft_deltas),
+        "total_latency_delta_s": sum(latency_deltas),
+        "mean_baseline_latency_s": _mean(
+            [
+                float(base["latency_s"])
+                for base, cand, _feature in pairs
+                if base.get("latency_s") is not None and cand.get("latency_s") is not None
+            ]
+        ),
+        "mean_candidate_latency_s": _mean(
+            [
+                float(cand["latency_s"])
+                for base, cand, _feature in pairs
+                if base.get("latency_s") is not None and cand.get("latency_s") is not None
+            ]
+        ),
+    }
+
+
+def _load_jsonl(path: Path) -> list[dict[str, Any]]:
+    rows: list[dict[str, Any]] = []
+    with path.open("r", encoding="utf-8") as handle:
+        for line in handle:
+            if line.strip():
+                rows.append(json.loads(line))
+    return rows
+
+
+def _numeric_values(rows: Iterable[dict[str, Any]], key: str) -> list[float]:
+    values: list[float] = []
+    for row in rows:
+        value = row.get(key)
+        if value is not None:
+            values.append(float(value))
+    return values
+
+
+def _stats(values: list[float]) -> dict[str, float] | None:
+    clean = [value for value in values if value is not None]
+    if not clean:
+        return None
+    clean.sort()
+    return {
+        "count": float(len(clean)),
+        "mean": statistics.fmean(clean),
+        "p50": _percentile(clean, 0.50),
+        "p90": _percentile(clean, 0.90),
+        "p99": _percentile(clean, 0.99),
+        "min": clean[0],
+        "max": clean[-1],
+    }
+
+
+def _percentile(sorted_values: list[float], percentile: float) -> float:
+    if len(sorted_values) == 1:
+        return sorted_values[0]
+    index = round((len(sorted_values) - 1) * percentile)
+    return sorted_values[index]
+
+
+def _mean(values: list[float]) -> float | None:
+    if not values:
+        return None
+    return statistics.fmean(values)
+
+
+def _token_bin(value: int | None) -> str:
+    if value is None:
+        return "none"
+    if value <= 0:
+        return "<=0"
+    if value <= 512:
+        return "1-512"
+    if value <= 2048:
+        return "513-2048"
+    if value <= 8192:
+        return "2049-8192"
+    if value <= 32768:
+        return "8193-32768"
+    return ">32768"

src/agentic_pd_hybrid/sampling.py

@@ -0,0 +1,295 @@
+from __future__ import annotations
+
+import hashlib
+import json
+from collections import defaultdict
+from dataclasses import asdict, dataclass
+from pathlib import Path
+from typing import Literal
+
+from agentic_pd_hybrid.trace import TraceRequest, load_trace
+
+
+SampleProfile = Literal["default", "small-append"]
+
+
+@dataclass(frozen=True)
+class SessionSampleConfig:
+    trace_path: Path
+    output_path: Path
+    target_duration_s: float = 600.0
+    start_time_s: float = 0.0
+    session_sample_rate: float = 1.0
+    min_turns: int = 1
+    max_requests: int | None = None
+    profile: SampleProfile = "default"
+    min_initial_input_tokens: int | None = None
+    max_initial_input_tokens: int | None = None
+    max_append_input_tokens: int | None = None
+    max_output_tokens: int | None = None
+    min_overlap_ratio: float | None = None
+
+
+@dataclass(frozen=True)
+class SessionSampleSummary:
+    input_trace_path: str
+    output_trace_path: str
+    request_count: int
+    session_count: int
+    multi_turn_session_count: int
+    start_time_s: float
+    end_time_s: float
+    sampled_duration_s: float
+    session_sample_rate: float
+    min_turns: int
+    profile: str
+    min_initial_input_tokens: int | None
+    max_initial_input_tokens: int | None
+    max_append_input_tokens: int | None
+    max_output_tokens: int | None
+    min_overlap_ratio: float | None
+    mean_append_input_tokens: float | None
+    mean_turn_overlap_ratio: float | None
+
+
+def sample_trace_sessions(config: SessionSampleConfig) -> SessionSampleSummary:
+    requests = load_trace(config.trace_path)
+    sessions: dict[str, list[TraceRequest]] = defaultdict(list)
+    for request in requests:
+        sessions[request.session_id].append(request)
+
+    filters = _resolve_filters(config)
+    eligible_sessions = {
+        session_id: session_requests
+        for session_id, session_requests in sessions.items()
+        if len(session_requests) >= filters.min_turns
+        and _session_matches_filters(session_requests, filters)
+        and _keep_session(session_id, config.session_sample_rate)
+    }
+    ordered_sessions = sorted(
+        eligible_sessions.values(),
+        key=lambda session_requests: session_requests[0].timestamp_s,
+    )
+
+    selected_requests: list[TraceRequest] = []
+    sampled_start: float | None = None
+    sampled_end: float | None = None
+    for session_requests in ordered_sessions:
+        session_first = session_requests[0].timestamp_s
+        if session_first < config.start_time_s:
+            continue
+
+        if sampled_start is None:
+            sampled_start = session_first
+
+        selected_requests.extend(session_requests)
+        sampled_end = max(request.timestamp_s for request in session_requests)
+
+        if config.max_requests is not None and len(selected_requests) >= config.max_requests:
+            break
+        if sampled_end - sampled_start >= config.target_duration_s:
+            break
+
+    selected_requests.sort(key=lambda request: request.timestamp_s)
+    if config.max_requests is not None:
+        selected_requests = selected_requests[: config.max_requests]
+
+    if not selected_requests:
+        raise ValueError("Sampling produced no requests; adjust the sampling arguments")
+
+    config.output_path.parent.mkdir(parents=True, exist_ok=True)
+    with config.output_path.open("w", encoding="utf-8") as handle:
+        for request in selected_requests:
+            payload = {
+                "request_id": request.request_id,
+                "session_id": request.session_id,
+                "chat_id": request.chat_id,
+                "parent_chat_id": request.parent_chat_id,
+                "timestamp": request.timestamp_s,
+                "input_length": request.input_length,
+                "output_length": request.output_length,
+                "type": request.request_type,
+                "turn": request.turn_id,
+                "hash_ids": list(request.hash_ids),
+            }
+            handle.write(json.dumps(payload, sort_keys=True) + "\n")
+
+    selected_session_ids = {request.session_id for request in selected_requests}
+    selected_session_requests = [
+        eligible_sessions[session_id] for session_id in selected_session_ids
+    ]
+    append_lengths = [
+        length
+        for session_requests in selected_session_requests
+        for length in _turn_append_lengths(session_requests)
+    ]
+    overlap_ratios = [
+        ratio
+        for session_requests in selected_session_requests
+        for ratio in _turn_overlap_ratios(session_requests)
+    ]
+    summary = SessionSampleSummary(
+        input_trace_path=str(config.trace_path),
+        output_trace_path=str(config.output_path),
+        request_count=len(selected_requests),
+        session_count=len(selected_session_ids),
+        multi_turn_session_count=sum(
+            1
+            for session_id in selected_session_ids
+            if len(eligible_sessions[session_id]) > 1
+        ),
+        start_time_s=selected_requests[0].timestamp_s,
+        end_time_s=selected_requests[-1].timestamp_s,
+        sampled_duration_s=selected_requests[-1].timestamp_s
+        - selected_requests[0].timestamp_s,
+        session_sample_rate=config.session_sample_rate,
+        min_turns=filters.min_turns,
+        profile=config.profile,
+        min_initial_input_tokens=filters.min_initial_input_tokens,
+        max_initial_input_tokens=filters.max_initial_input_tokens,
+        max_append_input_tokens=filters.max_append_input_tokens,
+        max_output_tokens=filters.max_output_tokens,
+        min_overlap_ratio=filters.min_overlap_ratio,
+        mean_append_input_tokens=_mean(append_lengths),
+        mean_turn_overlap_ratio=_mean(overlap_ratios),
+    )
+    summary_path = config.output_path.with_suffix(config.output_path.suffix + ".summary.json")
+    with summary_path.open("w", encoding="utf-8") as handle:
+        json.dump(asdict(summary), handle, indent=2, sort_keys=True)
+    return summary
+
+
+@dataclass(frozen=True)
+class _ResolvedFilters:
+    min_turns: int
+    min_initial_input_tokens: int | None
+    max_initial_input_tokens: int | None
+    max_append_input_tokens: int | None
+    max_output_tokens: int | None
+    min_overlap_ratio: float | None
+
+
+def _resolve_filters(config: SessionSampleConfig) -> _ResolvedFilters:
+    if config.profile == "default":
+        return _ResolvedFilters(
+            min_turns=config.min_turns,
+            min_initial_input_tokens=config.min_initial_input_tokens,
+            max_initial_input_tokens=config.max_initial_input_tokens,
+            max_append_input_tokens=config.max_append_input_tokens,
+            max_output_tokens=config.max_output_tokens,
+            min_overlap_ratio=config.min_overlap_ratio,
+        )
+
+    if config.profile != "small-append":
+        raise ValueError(f"Unsupported sample profile: {config.profile}")
+
+    return _ResolvedFilters(
+        min_turns=max(config.min_turns, 2),
+        min_initial_input_tokens=(
+            2048
+            if config.min_initial_input_tokens is None
+            else config.min_initial_input_tokens
+        ),
+        max_initial_input_tokens=(
+            16000
+            if config.max_initial_input_tokens is None
+            else config.max_initial_input_tokens
+        ),
+        max_append_input_tokens=(
+            2048
+            if config.max_append_input_tokens is None
+            else config.max_append_input_tokens
+        ),
+        max_output_tokens=(
+            2048 if config.max_output_tokens is None else config.max_output_tokens
+        ),
+        min_overlap_ratio=(
+            0.75 if config.min_overlap_ratio is None else config.min_overlap_ratio
+        ),
+    )
+
+
+def _session_matches_filters(
+    session_requests: list[TraceRequest],
+    filters: _ResolvedFilters,
+) -> bool:
+    ordered = sorted(
+        session_requests,
+        key=lambda request: (request.timestamp_s, request.turn_id, request.chat_id),
+    )
+    if not ordered:
+        return False
+
+    initial = ordered[0]
+    if (
+        filters.min_initial_input_tokens is not None
+        and initial.input_length < filters.min_initial_input_tokens
+    ):
+        return False
+    if (
+        filters.max_initial_input_tokens is not None
+        and initial.input_length > filters.max_initial_input_tokens
+    ):
+        return False
+    if filters.max_output_tokens is not None and any(
+        request.output_length > filters.max_output_tokens for request in ordered
+    ):
+        return False
+
+    append_lengths = _turn_append_lengths(ordered)
+    if filters.max_append_input_tokens is not None and any(
+        append_length <= 0 or append_length > filters.max_append_input_tokens
+        for append_length in append_lengths
+    ):
+        return False
+
+    overlap_ratios = _turn_overlap_ratios(ordered)
+    if filters.min_overlap_ratio is not None and any(
+        overlap_ratio < filters.min_overlap_ratio for overlap_ratio in overlap_ratios
+    ):
+        return False
+
+    return True
+
+
+def _turn_append_lengths(session_requests: list[TraceRequest]) -> list[int]:
+    ordered = sorted(
+        session_requests,
+        key=lambda request: (request.timestamp_s, request.turn_id, request.chat_id),
+    )
+    return [
+        current.input_length - (previous.input_length + previous.output_length)
+        for previous, current in zip(ordered, ordered[1:], strict=False)
+    ]
+
+
+def _turn_overlap_ratios(session_requests: list[TraceRequest]) -> list[float]:
+    ordered = sorted(
+        session_requests,
+        key=lambda request: (request.timestamp_s, request.turn_id, request.chat_id),
+    )
+    ratios: list[float] = []
+    for previous, current in zip(ordered, ordered[1:], strict=False):
+        if not current.hash_ids:
+            ratios.append(0.0)
+            continue
+        previous_blocks = set(previous.hash_ids)
+        overlap = sum(1 for block in current.hash_ids if block in previous_blocks)
+        ratios.append(overlap / len(current.hash_ids))
+    return ratios
+
+
+def _mean(values: list[int] | list[float]) -> float | None:
+    if not values:
+        return None
+    return sum(values) / len(values)
+
+
+def _keep_session(session_id: str, sample_rate: float) -> bool:
+    if sample_rate >= 1.0:
+        return True
+    if sample_rate <= 0.0:
+        return False
+    digest = hashlib.blake2b(session_id.encode("utf-8"), digest_size=8).digest()
+    bucket = int.from_bytes(digest, byteorder="big", signed=False) / 2**64
+    return bucket < sample_rate

src/agentic_pd_hybrid/stack.py

@@ -0,0 +1,229 @@
+from __future__ import annotations
+
+import os
+import signal
+import subprocess
+import time
+from dataclasses import dataclass
+from pathlib import Path
+
+import httpx
+
+from agentic_pd_hybrid.launcher import build_launch_plan
+from agentic_pd_hybrid.topology import SingleNodeTopology
+
+
+@dataclass
+class ManagedProcess:
+    name: str
+    command: tuple[str, ...]
+    process: subprocess.Popen[bytes]
+    log_path: Path
+
+
+@dataclass
+class ManagedPdStack:
+    topology: SingleNodeTopology
+    run_dir: Path
+    prefill_processes: list[ManagedProcess]
+    decode_processes: list[ManagedProcess]
+    direct_processes: list[ManagedProcess]
+    router_process: ManagedProcess | None
+
+    @property
+    def router_url(self) -> str:
+        return self.topology.router_url
+
+    def stop(self) -> None:
+        processes = (
+            ([self.router_process] if self.router_process is not None else [])
+            + self.direct_processes
+            + self.decode_processes
+            + self.prefill_processes
+        )
+        for managed in processes:
+            if managed.process.poll() is None:
+                os.killpg(os.getpgid(managed.process.pid), signal.SIGTERM)
+        deadline = time.time() + 20
+        for managed in processes:
+            if managed.process.poll() is not None:
+                continue
+            remaining = max(0.0, deadline - time.time())
+            try:
+                managed.process.wait(timeout=remaining)
+            except subprocess.TimeoutExpired:
+                if managed.process.poll() is None:
+                    os.killpg(os.getpgid(managed.process.pid), signal.SIGKILL)
+                    managed.process.wait(timeout=5)
+
+
+def launch_pd_stack(
+    *,
+    topology: SingleNodeTopology,
+    run_dir: Path,
+    prefill_policy: str,
+    decode_policy: str,
+    timeout_s: float = 1200.0,
+    router_request_timeout_s: float | None = None,
+    include_router: bool = True,
+    naive_dp: bool = False,
+) -> ManagedPdStack:
+    run_dir.mkdir(parents=True, exist_ok=True)
+    logs_dir = run_dir / "logs"
+    logs_dir.mkdir(parents=True, exist_ok=True)
+
+    plan = build_launch_plan(
+        topology,
+        prefill_policy=prefill_policy,
+        decode_policy=decode_policy,
+        include_router=include_router,
+        router_request_timeout_s=router_request_timeout_s,
+        naive_dp=naive_dp,
+    )
+
+    prefill_processes = [
+        _spawn_process(
+            name=f"prefill-{idx}",
+            command=command,
+            log_path=logs_dir / f"prefill-{idx}.log",
+            topology=topology,
+        )
+        for idx, command in enumerate(plan.prefill_commands)
+    ]
+    decode_processes = [
+        _spawn_process(
+            name=f"decode-{idx}",
+            command=command,
+            log_path=logs_dir / f"decode-{idx}.log",
+            topology=topology,
+        )
+        for idx, command in enumerate(plan.decode_commands)
+    ]
+    direct_processes = [
+        _spawn_process(
+            name=f"direct-{idx}",
+            command=command,
+            log_path=logs_dir / f"direct-{idx}.log",
+            topology=topology,
+        )
+        for idx, command in enumerate(plan.direct_commands)
+    ]
+
+    router_process: ManagedProcess | None = None
+    try:
+        for worker in topology.prefill_workers:
+            _wait_for_ready_endpoint(f"{worker.url}/v1/models", timeout_s=timeout_s)
+        for worker in topology.decode_workers:
+            _wait_for_ready_endpoint(f"{worker.url}/v1/models", timeout_s=timeout_s)
+        for worker in topology.direct_workers:
+            _wait_for_ready_endpoint(f"{worker.url}/v1/models", timeout_s=timeout_s)
+
+        if plan.router_command is not None:
+            router_process = _spawn_process(
+                name="router",
+                command=plan.router_command,
+                log_path=logs_dir / "router.log",
+                topology=topology,
+            )
+            _wait_for_ready_endpoint(f"{topology.router_url}/health", timeout_s=timeout_s)
+    except Exception:
+        stack = ManagedPdStack(
+            topology=topology,
+            run_dir=run_dir,
+            prefill_processes=prefill_processes,
+            decode_processes=decode_processes,
+            direct_processes=direct_processes,
+            router_process=router_process,
+        )
+        stack.stop()
+        raise
+
+    return ManagedPdStack(
+        topology=topology,
+        run_dir=run_dir,
+        prefill_processes=prefill_processes,
+        decode_processes=decode_processes,
+        direct_processes=direct_processes,
+        router_process=router_process,
+    )
+
+
+def _spawn_process(
+    *,
+    name: str,
+    command: tuple[str, ...],
+    log_path: Path,
+    topology: SingleNodeTopology,
+) -> ManagedProcess:
+    log_handle = log_path.open("wb")
+    env = _build_process_env(topology)
+    process = subprocess.Popen(
+        command,
+        stdout=log_handle,
+        stderr=subprocess.STDOUT,
+        env=env,
+        preexec_fn=os.setsid,
+    )
+    return ManagedProcess(
+        name=name,
+        command=command,
+        process=process,
+        log_path=log_path,
+    )
+
+
+def _build_process_env(topology: SingleNodeTopology) -> dict[str, str]:
+    env = os.environ.copy()
+    env["PYTHONDONTWRITEBYTECODE"] = "1"
+    env["PYTHONUNBUFFERED"] = "1"
+
+    # SGLang's PD bootstrap path uses `requests`; force localhost traffic to stay local.
+    for key in (
+        "http_proxy",
+        "https_proxy",
+        "all_proxy",
+        "HTTP_PROXY",
+        "HTTPS_PROXY",
+        "ALL_PROXY",
+    ):
+        env.pop(key, None)
+    env["NO_PROXY"] = "*"
+    env["no_proxy"] = "*"
+    env.setdefault("SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT", "600")
+    env.setdefault("SGLANG_DISAGGREGATION_WAITING_TIMEOUT", "600")
+    env.setdefault("FLASHINFER_DISABLE_VERSION_CHECK", "1")
+    if topology.force_rdma:
+        env["MOONCAKE_PROTOCOL"] = "rdma"
+        env["MC_MS_AUTO_DISC"] = "0"
+        if topology.ib_device:
+            env["MOONCAKE_DEVICE"] = topology.ib_device
+    elif topology.transfer_backend == "mooncake":
+        # Default to TCP when RDMA is not forced (e.g. loopback on same node)
+        env.setdefault("MOONCAKE_PROTOCOL", "tcp")
+
+    repo_root = Path(__file__).resolve().parents[2]
+    python_paths = [
+        str(repo_root / "src"),
+        str(repo_root / "third_party" / "sglang" / "python"),
+    ]
+    existing_pythonpath = env.get("PYTHONPATH")
+    if existing_pythonpath:
+        python_paths.append(existing_pythonpath)
+    env["PYTHONPATH"] = os.pathsep.join(python_paths)
+    return env
+
+
+def _wait_for_ready_endpoint(url: str, *, timeout_s: float) -> None:
+    start = time.perf_counter()
+    last_error: str | None = None
+    with httpx.Client(timeout=5.0, trust_env=False) as client:
+        while time.perf_counter() - start < timeout_s:
+            try:
+                response = client.get(url)
+                if response.status_code == 200:
+                    return
+                last_error = f"status={response.status_code}"
+            except Exception as exc:  # pragma: no cover
+                last_error = f"{type(exc).__name__}: {exc}"
+            time.sleep(1.0)
+    raise TimeoutError(f"Timed out waiting for {url} ({last_error})")

src/agentic_pd_hybrid/topology.py

@@ -0,0 +1,245 @@
+from __future__ import annotations
+
+from dataclasses import dataclass
+from pathlib import Path
+from typing import Literal
+
+
+WorkerRole = Literal["prefill", "decode", "direct"]
+
+
+@dataclass(frozen=True)
+class WorkerSpec:
+    role: WorkerRole
+    ordinal: int
+    gpu_ids: tuple[int, ...]
+    host: str
+    port: int
+    bootstrap_port: int | None = None
+
+    @property
+    def worker_id(self) -> str:
+        return f"{self.role}-{self.ordinal}"
+
+    @property
+    def url(self) -> str:
+        return f"http://{self.host}:{self.port}"
+
+    @property
+    def gpu_id(self) -> int:
+        return self.gpu_ids[0]
+
+    @property
+    def tp_size(self) -> int:
+        return len(self.gpu_ids)
+
+
+@dataclass(frozen=True)
+class SingleNodeTopology:
+    model_path: str
+    prefill_workers: tuple[WorkerSpec, ...]
+    decode_workers: tuple[WorkerSpec, ...]
+    direct_workers: tuple[WorkerSpec, ...]
+    router_host: str
+    router_port: int
+    transfer_backend: str
+    trust_remote_code: bool
+    force_rdma: bool = False
+    ib_device: str | None = None
+    extra_server_args: tuple[str, ...] = ()
+    prefill_extra_server_args: tuple[str, ...] = ()
+    decode_extra_server_args: tuple[str, ...] = ()
+    direct_extra_server_args: tuple[str, ...] = ()
+
+    @property
+    def model_name(self) -> str:
+        return Path(self.model_path).name
+
+    @property
+    def router_url(self) -> str:
+        return f"http://{self.router_host}:{self.router_port}"
+
+    @property
+    def route_workers(self) -> tuple[WorkerSpec, ...]:
+        if self.decode_workers:
+            return self.decode_workers
+        return self.direct_workers
+
+    def route_index(self, worker_id: str) -> int:
+        for idx, worker in enumerate(self.route_workers):
+            if worker.worker_id == worker_id:
+                return idx
+        raise KeyError(f"Unknown route worker: {worker_id}")
+
+
+def build_single_node_topology(
+    *,
+    model_path: str,
+    prefill_worker_count: int,
+    decode_worker_count: int,
+    direct_worker_count: int = 0,
+    prefill_tp_size: int = 1,
+    decode_tp_size: int = 1,
+    direct_tp_size: int = 1,
+    prefill_gpu_ids: tuple[int, ...] | None = None,
+    decode_gpu_ids: tuple[int, ...] | None = None,
+    direct_gpu_ids: tuple[int, ...] | None = None,
+    total_gpu_budget: int = 8,
+    host: str = "127.0.0.1",
+    router_port: int = 8000,
+    prefill_port_base: int = 30000,
+    decode_port_base: int = 31000,
+    direct_port_base: int = 32000,
+    bootstrap_port_base: int = 8998,
+    transfer_backend: str = "nixl",
+    force_rdma: bool = False,
+    trust_remote_code: bool = True,
+    ib_device: str | None = None,
+    extra_server_args: tuple[str, ...] = (),
+    prefill_extra_server_args: tuple[str, ...] = (),
+    decode_extra_server_args: tuple[str, ...] = (),
+    direct_extra_server_args: tuple[str, ...] = (),
+) -> SingleNodeTopology:
+    if prefill_worker_count < 0:
+        raise ValueError("prefill_worker_count must be >= 0")
+    if decode_worker_count < 0:
+        raise ValueError("decode_worker_count must be >= 0")
+    if direct_worker_count < 0:
+        raise ValueError("direct_worker_count must be >= 0")
+    if (
+        prefill_worker_count == 0
+        and decode_worker_count == 0
+        and direct_worker_count == 0
+    ):
+        raise ValueError("At least one worker must be configured")
+    if prefill_tp_size <= 0:
+        raise ValueError("prefill_tp_size must be >= 1")
+    if decode_tp_size <= 0:
+        raise ValueError("decode_tp_size must be >= 1")
+    if direct_tp_size <= 0:
+        raise ValueError("direct_tp_size must be >= 1")
+    if force_rdma and not ib_device:
+        raise ValueError("force_rdma requires --ib-device to be set")
+    if force_rdma and transfer_backend != "mooncake":
+        raise ValueError(
+            "force_rdma currently requires transfer_backend='mooncake' "
+            "to guarantee an RDMA path"
+        )
+
+    total_gpus_required = (
+        prefill_worker_count * prefill_tp_size
+        + decode_worker_count * decode_tp_size
+        + direct_worker_count * direct_tp_size
+    )
+    if total_gpus_required > total_gpu_budget:
+        raise ValueError(
+            "Single-node GPU budget exceeded: "
+            f"{prefill_worker_count} prefill x tp={prefill_tp_size} + "
+            f"{decode_worker_count} decode x tp={decode_tp_size} + "
+            f"{direct_worker_count} direct x tp={direct_tp_size} > "
+            f"{total_gpu_budget} GPUs"
+        )
+
+    if prefill_gpu_ids is None:
+        prefill_gpu_ids = tuple(range(prefill_worker_count * prefill_tp_size))
+    if decode_gpu_ids is None:
+        decode_gpu_ids = tuple(
+            range(
+                len(prefill_gpu_ids),
+                len(prefill_gpu_ids) + decode_worker_count * decode_tp_size,
+            )
+        )
+    if direct_gpu_ids is None:
+        direct_gpu_ids = tuple(
+            range(
+                len(prefill_gpu_ids) + len(decode_gpu_ids),
+                len(prefill_gpu_ids)
+                + len(decode_gpu_ids)
+                + direct_worker_count * direct_tp_size,
+            )
+        )
+
+    if len(prefill_gpu_ids) != prefill_worker_count * prefill_tp_size:
+        raise ValueError(
+            "prefill_gpu_ids length must equal prefill_worker_count * prefill_tp_size: "
+            f"{len(prefill_gpu_ids)} != {prefill_worker_count * prefill_tp_size}"
+        )
+    if len(decode_gpu_ids) != decode_worker_count * decode_tp_size:
+        raise ValueError(
+            "decode_gpu_ids length must equal decode_worker_count * decode_tp_size: "
+            f"{len(decode_gpu_ids)} != {decode_worker_count * decode_tp_size}"
+        )
+    if len(direct_gpu_ids) != direct_worker_count * direct_tp_size:
+        raise ValueError(
+            "direct_gpu_ids length must equal direct_worker_count * direct_tp_size: "
+            f"{len(direct_gpu_ids)} != {direct_worker_count * direct_tp_size}"
+        )
+    assigned_gpu_ids = prefill_gpu_ids + decode_gpu_ids + direct_gpu_ids
+    if len(set(assigned_gpu_ids)) != len(assigned_gpu_ids):
+        raise ValueError("prefill/decode/direct GPU IDs must be unique")
+    if any(gpu_id < 0 or gpu_id >= total_gpu_budget for gpu_id in assigned_gpu_ids):
+        raise ValueError(
+            "GPU IDs must fall within the single-node budget range "
+            f"[0, {total_gpu_budget - 1}]"
+        )
+
+    prefill_workers = tuple(
+        WorkerSpec(
+            role="prefill",
+            ordinal=idx,
+            gpu_ids=tuple(
+                prefill_gpu_ids[
+                    idx * prefill_tp_size : (idx + 1) * prefill_tp_size
+                ]
+            ),
+            host=host,
+            port=prefill_port_base + idx,
+            bootstrap_port=bootstrap_port_base + idx,
+        )
+        for idx in range(prefill_worker_count)
+    )
+    decode_workers = tuple(
+        WorkerSpec(
+            role="decode",
+            ordinal=idx,
+            gpu_ids=tuple(
+                decode_gpu_ids[
+                    idx * decode_tp_size : (idx + 1) * decode_tp_size
+                ]
+            ),
+            host=host,
+            port=decode_port_base + idx,
+        )
+        for idx in range(decode_worker_count)
+    )
+    direct_workers = tuple(
+        WorkerSpec(
+            role="direct",
+            ordinal=idx,
+            gpu_ids=tuple(
+                direct_gpu_ids[
+                    idx * direct_tp_size : (idx + 1) * direct_tp_size
+                ]
+            ),
+            host=host,
+            port=direct_port_base + idx,
+        )
+        for idx in range(direct_worker_count)
+    )
+
+    return SingleNodeTopology(
+        model_path=model_path,
+        prefill_workers=prefill_workers,
+        decode_workers=decode_workers,
+        direct_workers=direct_workers,
+        router_host=host,
+        router_port=router_port,
+        transfer_backend=transfer_backend,
+        trust_remote_code=trust_remote_code,
+        force_rdma=force_rdma,
+        ib_device=ib_device,
+        extra_server_args=extra_server_args,
+        prefill_extra_server_args=prefill_extra_server_args,
+        decode_extra_server_args=decode_extra_server_args,
+        direct_extra_server_args=direct_extra_server_args,
+    )

src/agentic_pd_hybrid/trace.py

@@ -0,0 +1,106 @@
+from __future__ import annotations
+
+import json
+from dataclasses import dataclass
+from pathlib import Path
+
+
+@dataclass(frozen=True)
+class TraceRequest:
+    request_id: str
+    session_id: str
+    chat_id: int
+    parent_chat_id: int
+    timestamp_s: float
+    input_length: int
+    output_length: int
+    request_type: str
+    turn_id: int
+    hash_ids: tuple[int, ...]
+
+
+def load_trace(path: Path, *, request_limit: int | None = None) -> list[TraceRequest]:
+    chat_to_session: dict[int, str] = {}
+    requests: list[TraceRequest] = []
+
+    with path.open("r", encoding="utf-8") as handle:
+        for index, line in enumerate(handle):
+            if request_limit is not None and len(requests) >= request_limit:
+                break
+
+            payload = json.loads(line)
+            chat_id = int(payload["chat_id"])
+            parent_chat_id = int(payload["parent_chat_id"])
+            session_id = _resolve_session_id(
+                chat_id=chat_id,
+                parent_chat_id=parent_chat_id,
+                chat_to_session=chat_to_session,
+            )
+            turn_id = int(payload["turn"])
+            request_id = f"{session_id}:{turn_id}:{chat_id}:{index}"
+            requests.append(
+                TraceRequest(
+                    request_id=request_id,
+                    session_id=session_id,
+                    chat_id=chat_id,
+                    parent_chat_id=parent_chat_id,
+                    timestamp_s=float(payload["timestamp"]),
+                    input_length=int(payload["input_length"]),
+                    output_length=int(payload["output_length"]),
+                    request_type=str(payload["type"]),
+                    turn_id=turn_id,
+                    hash_ids=tuple(int(item) for item in payload.get("hash_ids", [])),
+                )
+            )
+
+    return requests
+
+
+def build_synthetic_prompt(
+    request: TraceRequest,
+    *,
+    block_token_budget: int = 24,
+) -> str:
+    return " ".join(build_synthetic_prompt_tokens(request, block_token_budget=block_token_budget))
+
+
+def build_synthetic_prompt_tokens(
+    request: TraceRequest,
+    *,
+    block_token_budget: int = 24,
+) -> list[str]:
+    tokens: list[str] = []
+    for hash_id in request.hash_ids:
+        for offset in range(block_token_budget):
+            tokens.append(f"blk{hash_id}_{offset}")
+
+    while len(tokens) < request.input_length:
+        tokens.append(f"fill_{len(tokens) % 64}")
+
+    return tokens[: request.input_length]
+
+
+def build_synthetic_append_chunk(
+    request: TraceRequest,
+    append_length: int,
+) -> str:
+    if append_length <= 0:
+        return ""
+    return " ".join(
+        f"turn{request.turn_id}_append_{request.chat_id}_{offset}"
+        for offset in range(append_length)
+    )
+
+
+def _resolve_session_id(
+    *,
+    chat_id: int,
+    parent_chat_id: int,
+    chat_to_session: dict[int, str],
+) -> str:
+    if parent_chat_id < 0:
+        session_id = str(chat_id)
+    else:
+        session_id = chat_to_session.get(parent_chat_id, str(parent_chat_id))
+    chat_to_session[chat_id] = session_id
+    return session_id

src/agentic_pd_hybrid/trace_profiles.py

@@ -0,0 +1,127 @@
+from __future__ import annotations
+
+import json
+from collections import defaultdict
+from dataclasses import asdict, dataclass
+from math import ceil
+from pathlib import Path
+
+from agentic_pd_hybrid.trace import TraceRequest, load_trace
+
+
+BLOCK_TOKEN_BUDGET = 24
+
+
+@dataclass(frozen=True)
+class NormalizeTraceLengthsConfig:
+    trace_path: Path
+    output_path: Path
+    initial_input_length: int = 10_000
+    append_input_length: int = 1_000
+    output_length: int = 1_000
+    max_requests: int | None = None
+
+
+@dataclass(frozen=True)
+class NormalizeTraceLengthsSummary:
+    input_trace_path: str
+    output_trace_path: str
+    request_count: int
+    session_count: int
+    multi_turn_session_count: int
+    initial_input_length: int
+    append_input_length: int
+    output_length: int
+    max_turns_per_session: int
+    max_input_length: int
+
+
+def normalize_trace_lengths(
+    config: NormalizeTraceLengthsConfig,
+) -> NormalizeTraceLengthsSummary:
+    if config.initial_input_length < 0:
+        raise ValueError("initial_input_length must be >= 0")
+    if config.append_input_length < 0:
+        raise ValueError("append_input_length must be >= 0")
+    if config.output_length < 0:
+        raise ValueError("output_length must be >= 0")
+
+    requests = load_trace(config.trace_path, request_limit=config.max_requests)
+    sessions: dict[str, list[TraceRequest]] = defaultdict(list)
+    for request in requests:
+        sessions[request.session_id].append(request)
+
+    normalized_records: list[dict[str, object]] = []
+    max_turns_per_session = 0
+    max_input_length = 0
+
+    for session_idx, session_id in enumerate(sorted(sessions, key=_session_sort_key)):
+        session_requests = sorted(
+            sessions[session_id],
+            key=lambda request: (request.timestamp_s, request.turn_id, request.chat_id),
+        )
+        max_turns_per_session = max(max_turns_per_session, len(session_requests))
+        base_block_count = ceil(config.initial_input_length / BLOCK_TOKEN_BUDGET)
+        base_hash_ids = [
+            _hash_id_for(session_idx=session_idx, block_idx=block_idx)
+            for block_idx in range(base_block_count)
+        ]
+
+        for turn_idx, request in enumerate(session_requests):
+            input_length = config.initial_input_length + turn_idx * (
+                config.append_input_length + config.output_length
+            )
+            total_block_count = ceil(input_length / BLOCK_TOKEN_BUDGET)
+            hash_ids = base_hash_ids + [
+                _hash_id_for(
+                    session_idx=session_idx,
+                    block_idx=base_block_count + append_block_idx,
+                )
+                for append_block_idx in range(max(0, total_block_count - base_block_count))
+            ]
+            max_input_length = max(max_input_length, input_length)
+            normalized_records.append(
+                {
+                    "chat_id": request.chat_id,
+                    "parent_chat_id": request.parent_chat_id,
+                    "timestamp": request.timestamp_s,
+                    "input_length": input_length,
+                    "output_length": config.output_length,
+                    "type": request.request_type,
+                    "turn": request.turn_id,
+                    "hash_ids": hash_ids,
+                }
+            )
+
+    normalized_records.sort(key=lambda item: float(item["timestamp"]))
+    config.output_path.parent.mkdir(parents=True, exist_ok=True)
+    with config.output_path.open("w", encoding="utf-8") as handle:
+        for record in normalized_records:
+            handle.write(json.dumps(record, sort_keys=True) + "\n")
+
+    summary = NormalizeTraceLengthsSummary(
+        input_trace_path=str(config.trace_path),
+        output_trace_path=str(config.output_path),
+        request_count=len(normalized_records),
+        session_count=len(sessions),
+        multi_turn_session_count=sum(
+            1 for session_requests in sessions.values() if len(session_requests) > 1
+        ),
+        initial_input_length=config.initial_input_length,
+        append_input_length=config.append_input_length,
+        output_length=config.output_length,
+        max_turns_per_session=max_turns_per_session,
+        max_input_length=max_input_length,
+    )
+    summary_path = config.output_path.with_suffix(config.output_path.suffix + ".summary.json")
+    with summary_path.open("w", encoding="utf-8") as handle:
+        json.dump(asdict(summary), handle, indent=2, sort_keys=True)
+    return summary
+
+
+def _hash_id_for(*, session_idx: int, block_idx: int) -> int:
+    return session_idx * 1_000_000 + block_idx
+
+
+def _session_sort_key(session_id: str) -> tuple[int, str]:
+    return (0, session_id) if session_id.isdigit() else (1, session_id)

third_party/sglang/.claude/skills/add-jit-kernel/SKILL.md

@@ -0,0 +1,607 @@
+---
+name: add-jit-kernel
+description: Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang's jit_kernel module
+---
+
+# Tutorial: Adding a New JIT Kernel to SGLang
+
+This tutorial walks through adding a simple element-wise scale operation as a JIT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.
+
+## Goal
+
+Add a new operation that scales each element of a tensor by a scalar factor:
+
+- Input: tensor `x` (CUDA) and scalar `factor` (float, passed at runtime)
+- Output: `x * factor` (element-wise), allocated internally
+- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**
+
+## When to use JIT vs AOT (`sgl-kernel`)
+
+- **JIT (`jit_kernel`)**: prefer this first for kernels that do **not** depend on CUTLASS or another large C++ project. It is the default choice for lightweight kernels that benefit from rapid iteration and first-use compilation.
+- **AOT (`sgl-kernel`)**: prefer this when the kernel **does** depend on CUTLASS or another large C++ project, or when it should live in `sgl-kernel/` and participate in the wheel build / torch op registration flow.
+- **Exception**: kernels that depend on `flashinfer`, or on CUTLASS that is already provided through `flashinfer`, can still be implemented as `jit_kernel`.
+
+---
+
+## Common Abstractions in `python/sglang/jit_kernel/include/sgl_kernel/`
+
+**Always prefer these abstractions over raw CUDA primitives.** They provide safety, readability, and consistency with the rest of the codebase.
+
+**Important include rule:** for every `#include <sgl_kernel/...>` line, add a short trailing comment explaining why that header is included (for example `// For TensorMatcher, SymbolicSize, SymbolicDevice`). This matches the current JIT kernel style and keeps include usage self-documenting.
+
+### `utils.h` — Host-side utilities
+
+```cpp
+#include <sgl_kernel/utils.h>
+```
+
+- **`host::RuntimeCheck(cond, args...)`** — Assert a condition at runtime; throws `PanicError` with file/line info on failure. Prefer this over bare `assert`.
+- **`host::Panic(args...)`** — Unconditionally throw a `PanicError` with a descriptive message.
+- **`host::div_ceil(a, b)`** — Integer ceiling division `(a + b - 1) / b`.
+- **`host::irange(n)`** / **`host::irange(start, end)`** — Range views for cleaner loops.
+- **`host::pointer::offset(ptr, offsets...)`** — Byte-safe pointer arithmetic on `void*`. Use this instead of raw casts.
+
+### `utils.cuh` — Device-side utilities + `LaunchKernel`
+
+```cpp
+#include <sgl_kernel/utils.cuh>
+```
+
+- **Type aliases**: `fp16_t`, `bf16_t`, `fp32_t`, `fp8_e4m3_t`, `fp8_e5m2_t` and their packed variants `fp16x2_t`, `bf16x2_t`, `fp32x2_t`, etc.
+- **`SGL_DEVICE`** — Expands to `__forceinline__ __device__`. Use on all device functions.
+- **`device::kWarpThreads`** — Constant `32`.
+- **`device::load_as<T>(ptr, offset)`** / **`device::store_as<T>(ptr, val, offset)`** — Type-safe loads/stores from `void*`.
+- **`device::pointer::offset(ptr, offsets...)`** — Pointer arithmetic on device.
+- **`host::LaunchKernel(grid, block, device_or_stream [, smem])`** — RAII kernel launcher that:
+  - Resolves the CUDA stream from a `DLDevice` via TVM-FFI automatically.
+  - Checks the CUDA error with file/line info after launch via `operator()(kernel, args...)`.
+  - Supports `.enable_pdl(bool)` for PDL (Programmatic Dependent Launch, SM90+).
+- **`host::RuntimeDeviceCheck(cudaError_t)`** — Check a CUDA error; throw on failure.
+
+### `tensor.h` — Tensor validation (`TensorMatcher`, Symbolic types)
+
+```cpp
+#include <sgl_kernel/tensor.h>
+```
+
+This is the **primary validation API** for all kernel launchers. Use it to validate every `tvm::ffi::TensorView` argument.
+
+- **`host::SymbolicSize{"name"}`** — A named symbolic dimension. Call `.set_value(n)` to pin it, `.unwrap()` to extract after verification.
+- **`host::SymbolicDType`** — Symbolic dtype. Use `.set_options<Ts...>()` to restrict allowed types.
+- **`host::SymbolicDevice`** — Symbolic device. Use `.set_options<kDLCUDA>()` to restrict to CUDA.
+- **`host::TensorMatcher({dims...})`** — Fluent builder for tensor validation:
+  - `.with_dtype<T>()` — require a specific C++ type (e.g. `fp16_t`)
+  - `.with_dtype<T1, T2, ...>()` — allow a set of types
+  - `.with_device<kDLCUDA>(device_sym)` — require CUDA and bind the checked device to a `SymbolicDevice`
+  - `.with_strides({strides...})` — validate strides (omit to require contiguous)
+  - `.verify(tensor_view)` — execute the check; throws `PanicError` with full context on failure; **chainable** (`verify(a).verify(b)` to check multiple tensors with the same shape)
+
+**Typical pattern:**
+```cpp
+auto N = SymbolicSize{"num_elements"};
+auto device = SymbolicDevice{};
+device.set_options<kDLCUDA>();
+TensorMatcher({N})  //
+    .with_dtype<fp16_t>()
+    .with_device<kDLCUDA>(device)
+    .verify(dst)
+    .verify(src);  // same shape, dtype, device as dst
+const size_t n = N.unwrap();
+const DLDevice dev = device.unwrap();
+```
+
+### `type.cuh` — `dtype_trait<T>` and `packed_t<T>`
+
+```cpp
+#include <sgl_kernel/type.cuh>
+```
+
+- **`dtype_trait<T>`** — Static trait struct for each scalar type. Provides:
+  - `dtype_trait<T>::from(value)` — convert from another type (e.g. `fp32_t` → `fp16_t`)
+  - `dtype_trait<T>::abs/sqrt/rsqrt/exp/sin/cos(x)` — type-dispatched unary math (primarily for `fp32_t`)
+  - `dtype_trait<T>::max/min(x, y)` — type-dispatched binary math (primarily for `fp32_t`)
+- **`packed_t<T>`** — Two-element packed alias: `packed_t<fp16_t>` = `fp16x2_t`, `packed_t<bf16_t>` = `bf16x2_t`, `packed_t<fp32_t>` = `fp32x2_t`. Use for vectorized loads/stores.
+- **`device::cast<To, From>(value)`** — Type-safe cast using `dtype_trait`, e.g. `cast<fp32x2_t, fp16x2_t>(v)`.
+
+### `vec.cuh` — Vectorized memory access (`AlignedVector`)
+
+```cpp
+#include <sgl_kernel/vec.cuh>
+```
+
+- **`device::AlignedVector<T, N>`** — Aligned storage for N elements of type T. N must be a power of two, `sizeof(T)*N <= 32`. Enables vectorized loads/stores for bandwidth efficiency. In terms of API/codegen constraints, the upper bound is 256-bit; in practice, 128-bit is the portable default, while 256-bit vectorization is typically only viable on `SM100+` and should be gated by an architecture check when needed.
+  - `.load(ptr, offset)` — vectorized load from `ptr[offset]`
+  - `.store(ptr, offset)` — vectorized store to `ptr[offset]`
+  - `.fill(value)` — fill all lanes
+  - `operator[](i)` — element access
+
+### `tile.cuh` — `tile::Memory` (strided memory access pattern)
+
+```cpp
+#include <sgl_kernel/tile.cuh>
+```
+
+- `tile::Memory<T>` is fundamentally a **1D cooperative accessor** over a contiguous region.
+- **`device::tile::Memory<T>::cta(blockDim.x)`** — Creates a tile accessor where each thread handles `tid = threadIdx.x` with stride `tsize` (for `cta(blockDim.x)`, this is `blockDim.x`). Common for loops over a 1D array.
+- **`.load(ptr, offset)`** — loads `ptr[tid + offset * tsize]`
+- **`.store(ptr, val, offset)`** — stores to `ptr[tid + offset * tsize]`
+- **`.in_bound(n, offset)`** — boundary check
+
+For a **2D tile**, either flatten `(row, col)` into a linear tile index first, or compute the address manually with `ptr[row * stride + col]` using your thread/block coordinates.
+
+### `math.cuh` — Device math (`device::math::`)
+
+```cpp
+#include <sgl_kernel/math.cuh>
+```
+
+- `device::math::max/min<T>(a, b)` — type-dispatched binary math via `dtype_trait`
+- `device::math::abs/sqrt/rsqrt/exp/sin/cos<T>(x)` — type-dispatched unary math via `dtype_trait`
+
+### `warp.cuh` — Warp-level primitives
+
+```cpp
+#include <sgl_kernel/warp.cuh>
+```
+
+- `device::warp::reduce_sum<T>(value)` — warp-level sum reduction via `__shfl_xor_sync`
+- `device::warp::reduce_max<T>(value)` — warp-level max reduction
+
+### `cta.cuh` — CTA-level primitives
+
+```cpp
+#include <sgl_kernel/cta.cuh>
+```
+
+- `device::cta::reduce_max<T>(value, smem, min_value)` — CTA-wide max using shared memory + warp reduction. Caller is responsible for a `__syncthreads()` after if the result in `smem[0]` is needed.
+
+### `atomic.cuh` — Atomic operations
+
+```cpp
+#include <sgl_kernel/atomic.cuh>
+```
+
+- `device::atomic::max(float* addr, float value)` — float atomic max (handles negative values correctly via bit tricks).
+
+### `runtime.cuh` — Occupancy and device info
+
+```cpp
+#include <sgl_kernel/runtime.cuh>
+```
+
+- `host::runtime::get_blocks_per_sm(kernel, block_dim)` — max active blocks per SM (occupancy)
+- `host::runtime::get_sm_count(device_id)` — number of SMs on the device
+- `host::runtime::get_cc_major(device_id)` — compute capability major version
+
+**Persistent kernel pattern** (cap blocks to SM count × occupancy):
+```cpp
+static const uint32_t max_occ = runtime::get_blocks_per_sm(kernel, kBlockSize);
+static const uint32_t num_sm  = runtime::get_sm_count(device.unwrap().device_id);
+const auto num_blocks = std::min(num_sm * max_occ, div_ceil(n, kBlockSize));
+LaunchKernel(num_blocks, kBlockSize, device.unwrap())(kernel, params);
+```
+
+---
+
+## Step 0 (optional): Generate a `.clangd` config for better IDE support
+
+```bash
+python -m sglang.jit_kernel
+```
+
+---
+
+## Step 1: Implement the CUDA kernel in `jit_kernel/csrc/`
+
+Create `python/sglang/jit_kernel/csrc/elementwise/scale.cuh`.
+
+The implementation fully uses the project abstractions described above:
+
+```cpp
+#include <sgl_kernel/tensor.h>   // For TensorMatcher, SymbolicSize, SymbolicDevice
+#include <sgl_kernel/type.cuh>   // For dtype_trait, fp16_t, bf16_t, fp32_t
+#include <sgl_kernel/utils.h>    // For RuntimeCheck, div_ceil
+#include <sgl_kernel/utils.cuh>  // For LaunchKernel, SGL_DEVICE
+#include <sgl_kernel/vec.cuh>    // For AlignedVector
+
+#include <dlpack/dlpack.h>
+#include <tvm/ffi/container/tensor.h>
+
+namespace {
+
+// ----------------------------------------------------------------
+// Kernel: element-wise scale using vectorized 128-bit loads/stores
+// T       = fp16_t | bf16_t | fp32_t
+// kVecN   = number of elements per vector load (e.g. 8 for fp16)
+// factor  = runtime scale factor
+// ----------------------------------------------------------------
+template <typename T, int kVecN>
+__global__ void scale_kernel(T* __restrict__ dst,
+                              const T* __restrict__ src,
+                              float factor,
+                              uint32_t n_total) {
+  using vec_t = device::AlignedVector<T, kVecN>;
+  const uint32_t n_vecs = n_total / kVecN;
+
+  // --- vectorised body ---
+  const uint32_t vec_stride = blockDim.x * gridDim.x;
+  for (uint32_t vi = blockIdx.x * blockDim.x + threadIdx.x;
+       vi < n_vecs;
+       vi += vec_stride) {
+    vec_t v;
+    v.load(src, vi);
+#pragma unroll
+    for (int i = 0; i < kVecN; ++i) {
+      v[i] = static_cast<T>(static_cast<float>(v[i]) * factor);
+    }
+    v.store(dst, vi);
+  }
+
+  // --- scalar tail ---
+  const uint32_t base = n_vecs * kVecN;
+  const uint32_t scalar_stride = blockDim.x * gridDim.x;
+  for (uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;
+       base + i < n_total;
+       i += scalar_stride) {
+    dst[base + i] = static_cast<T>(static_cast<float>(src[base + i]) * factor);
+  }
+}
+
+// ----------------------------------------------------------------
+// Launcher: validates tensors, selects vector width, launches kernel
+// ----------------------------------------------------------------
+template <typename T>
+void scale(tvm::ffi::TensorView dst, tvm::ffi::TensorView src, float factor) {
+  using namespace host;
+
+  // 1. Validate input tensors with TensorMatcher
+  SymbolicSize N = {"num_elements"};
+  SymbolicDevice device_;
+  device_.set_options<kDLCUDA>();
+
+  TensorMatcher({N})  //
+      .with_dtype<T>()
+      .with_device<kDLCUDA>(device_)
+      .verify(dst)
+      .verify(src);  // same shape / dtype / device as dst
+
+  const uint32_t n = static_cast<uint32_t>(N.unwrap());
+  const DLDevice device = device_.unwrap();
+
+  RuntimeCheck(n > 0, "scale: num_elements must be > 0, got ", n);
+
+  // 2. Choose vector width for 128-bit loads (16 bytes)
+  //    fp16/bf16: 8 elements × 2 bytes = 16 bytes
+  //    fp32:      4 elements × 4 bytes = 16 bytes
+  constexpr int kVecN = 16 / sizeof(T);
+  const uint32_t n_work_items = div_ceil(n, static_cast<uint32_t>(kVecN));
+
+  // 3. Launch
+  constexpr uint32_t kBlockSize = 256;
+  const uint32_t grid = div_ceil(n_work_items, kBlockSize);
+
+  LaunchKernel(grid, kBlockSize, device)(
+      scale_kernel<T, kVecN>,
+      static_cast<T*>(dst.data_ptr()),
+      static_cast<const T*>(src.data_ptr()),
+      factor,
+      n);
+}
+
+}  // namespace
+```
+
+**Key points:**
+
+- Include headers from `sgl_kernel/` — **not** raw CUDA headers for anything already covered
+- Add a short trailing `// For ...` explanation to every `#include <sgl_kernel/...>` line
+- Use `TensorMatcher` for all tensor validation; never manually check shape/dtype/device
+- Use `AlignedVector` for vectorised 128-bit loads/stores — significant bandwidth win
+- Use `LaunchKernel` — it resolves the stream and checks errors automatically
+- Use `RuntimeCheck` for runtime assertions with useful error messages
+- Prefer passing runtime scalars like `factor` directly unless compile-time specialisation is genuinely required
+- `fp16_t` / `bf16_t` / `fp32_t` are the project's type aliases (from `utils.cuh`)
+- `device::cast<To, From>` or `dtype_trait<T>::from(val)` for cross-type conversions
+- `device::math::` functions for device math instead of bare `__` intrinsics
+
+---
+
+## Step 2: Add the Python wrapper in `jit_kernel/`
+
+Create `python/sglang/jit_kernel/scale.py`:
+
+```python
+from __future__ import annotations
+
+from typing import TYPE_CHECKING
+
+import torch
+
+from sglang.jit_kernel.utils import cache_once, load_jit, make_cpp_args
+
+if TYPE_CHECKING:
+    from tvm_ffi.module import Module
+
+
+@cache_once
+def _jit_scale_module(dtype: torch.dtype) -> Module:
+    """Compile and cache the JIT scale module for a given dtype."""
+    args = make_cpp_args(dtype)
+    return load_jit(
+        "scale",
+        *args,
+        cuda_files=["elementwise/scale.cuh"],
+        cuda_wrappers=[("scale", f"scale<{args}>")],
+    )
+
+
+def scale(src: torch.Tensor, factor: float, out: torch.Tensor | None = None) -> torch.Tensor:
+    """
+    Element-wise scale: dst = src * factor.
+
+    Supported dtypes: torch.float16, torch.bfloat16, torch.float32.
+
+    Parameters
+    ----------
+    src    : CUDA tensor (FP16 / BF16 / FP32)
+    factor : scale factor
+    out    : optional pre-allocated output tensor (same shape/dtype as src)
+
+    Returns
+    -------
+    Scaled tensor (dst = src * factor).
+    """
+    if not src.is_cuda:
+        raise RuntimeError("src must be a CUDA tensor")
+    if src.dtype not in (torch.float16, torch.bfloat16, torch.float32):
+        raise RuntimeError(
+            f"Unsupported dtype {src.dtype}. Supported: float16, bfloat16, float32"
+        )
+    if out is None:
+        out = torch.empty_like(src)
+    else:
+        if out.shape != src.shape:
+            raise RuntimeError("out shape must match src")
+        if out.dtype != src.dtype:
+            raise RuntimeError("out dtype must match src")
+        if out.device != src.device:
+            raise RuntimeError("out device must match src")
+
+    # Keep the Python wrapper thin, but still enforce the basic preconditions
+    # that the current JIT/FFI path does not reject safely on its own.
+    module = _jit_scale_module(src.dtype)
+    module.scale(out, src, factor)
+    return out
+```
+
+**Key points:**
+
+- Use `cache_once` — **not** `functools.lru_cache` (incompatible with `torch.compile`)
+- `load_jit` first arg(s) form the unique build marker; same marker = same cached binary
+- Only include compile-time specialisation knobs in the build marker; runtime values like `factor` should stay runtime unless the kernel truly needs templating
+- `cuda_wrappers`: `(export_name, kernel_symbol)` — `export_name` is called from Python
+- `make_cpp_args(dtype, ...)` converts `torch.dtype` to C++ type alias:
+- Keep Python launchers thin, but still validate the basic invariants (`is_cuda`, supported dtype, `out` metadata). In the current JIT/FFI path, invalid tensors are not always rejected safely before launch
+
+| `torch.dtype`      | C++ type   |
+|--------------------|------------|
+| `torch.float16`    | `fp16_t`   |
+| `torch.bfloat16`   | `bf16_t`   |
+| `torch.float32`    | `fp32_t`   |
+
+---
+
+## Step 3 (optional): Tune JIT build flags
+
+```python
+return load_jit(
+    "scale",
+    *args,
+    cuda_files=["elementwise/scale.cuh"],
+    cuda_wrappers=[("scale", f"scale<{args}>")],
+    extra_cuda_cflags=["-O3", "--use_fast_math"],
+)
+```
+
+If your kernel requires SM90+, raise a clear Python error before calling `load_jit`:
+
+```python
+if torch.cuda.get_device_capability()[0] < 9:
+    raise RuntimeError("This kernel requires SM90 (Hopper) or later")
+```
+
+---
+
+## Step 4: Write tests (required)
+
+JIT kernel tests live under `python/sglang/jit_kernel/tests/`. **CI does not run `pytest` in that directory directly.** The unified runner `test/run_suite.py` discovers every `test_*.py` there (and every `bench_*.py` under `benchmark/`), collects `register_*_ci(...)` calls by **statically parsing each file’s AST**, and executes the selected suite. Every test file must register at least one CUDA entry or the collector fails its sanity check.
+
+- **PR / per-commit CUDA suites** (see `test/run_suite.py` → `PER_COMMIT_SUITES`): JIT unit tests use `stage-b-kernel-unit-1-gpu-large` (see `.github/workflows/pr-test-jit-kernel.yml`: `python3 run_suite.py --hw cuda --suite stage-b-kernel-unit-1-gpu-large`).
+- **Nightly kernel suite**: `nightly-kernel-1-gpu` with `--nightly` — typically used with `SGLANG_JIT_KERNEL_RUN_FULL_TESTS=1` in CI for expanded parameter grids (see `python/sglang/jit_kernel/utils.py` → `should_run_full_tests` / `get_ci_test_range`). Wired in `.github/workflows/nightly-test-nvidia.yml` (e.g. `python3 run_suite.py --hw cuda --suite nightly-kernel-1-gpu --nightly --continue-on-error`).
+
+Registration pattern (module level, **literal** `est_time` and `suite` strings — required for AST parsing):
+
+```python
+from sglang.test.ci.ci_register import register_cuda_ci
+
+register_cuda_ci(est_time=30, suite="stage-b-kernel-unit-1-gpu-large")
+# Optional second registration: same file also listed under the nightly kernel suite
+# register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True)
+```
+
+Keep `est_time` and `suite` as literal values. `run_suite.py` collects them from the file AST, so computed values and helper wrappers can break CI discovery.
+
+Use `register_cuda_ci(..., disabled="reason")` if the file must stay in-tree but should be skipped in CI (e.g. multi-GPU only).
+
+**Run like CI** (from repo root):
+
+```bash
+cd test && python3 run_suite.py --hw cuda --suite stage-b-kernel-unit-1-gpu-large
+```
+
+For fast iteration you can still run `pytest` on a single file locally; CI coverage is via `run_suite.py`.
+
+Create `python/sglang/jit_kernel/tests/test_scale.py`:
+
+```python
+import pytest
+import torch
+from sglang.jit_kernel.scale import scale
+from sglang.test.ci.ci_register import register_cuda_ci
+
+register_cuda_ci(est_time=30, suite="stage-b-kernel-unit-1-gpu-large")
+
+
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float32])
+@pytest.mark.parametrize("size", [1, 127, 128, 1024, 4097])  # cover tail remainder
+@pytest.mark.parametrize("factor", [0.5, 1.0, 2.0, 3.0])
+def test_scale_correctness(dtype, size, factor):
+    src = torch.randn(size, dtype=dtype, device="cuda")
+    out = scale(src, factor)
+    expected = src * factor
+
+    rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)
+    torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)
+
+
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float32])
+def test_scale_out_param(dtype):
+    src = torch.randn(1024, dtype=dtype, device="cuda")
+    out = torch.empty_like(src)
+    result = scale(src, 2.0, out=out)
+    assert result is out
+    torch.testing.assert_close(out, src * 2.0, rtol=1e-2, atol=1e-2)
+
+
+def test_scale_cpu_error():
+    src = torch.randn(128, dtype=torch.float16)  # CPU tensor
+    with pytest.raises(RuntimeError, match="CUDA"):
+        scale(src, 2.0)
+
+
+def test_scale_unsupported_dtype():
+    src = torch.randint(0, 10, (128,), dtype=torch.int32, device="cuda")
+    with pytest.raises(RuntimeError, match="dtype"):
+        scale(src, 2.0)
+
+
+if __name__ == "__main__":
+    import sys
+    sys.exit(pytest.main([__file__, "-v", "-s"]))
+```
+
+---
+
+## Step 5: Add a benchmark (required)
+
+Benchmarks are `bench_*.py` files under `python/sglang/jit_kernel/benchmark/`. They are picked up by the same `run_suite.py` machinery as unit tests. Register them for **`stage-b-kernel-benchmark-1-gpu-large`** (PR JIT benchmark job: `python3 run_suite.py --hw cuda --suite stage-b-kernel-benchmark-1-gpu-large`).
+
+Create `python/sglang/jit_kernel/benchmark/bench_scale.py`:
+
+```python
+import itertools
+
+import torch
+import triton
+import triton.testing
+
+from sglang.jit_kernel.benchmark.utils import (
+    DEFAULT_DEVICE,
+    DEFAULT_DTYPE,
+    get_benchmark_range,
+    run_benchmark,
+)
+from sglang.jit_kernel.scale import scale as jit_scale
+from sglang.test.ci.ci_register import register_cuda_ci
+
+register_cuda_ci(est_time=6, suite="stage-b-kernel-benchmark-1-gpu-large")
+
+SIZE_LIST = get_benchmark_range(
+    full_range=[2**n for n in range(10, 20)],  # 1K … 512K elements
+    ci_range=[4096, 65536],
+)
+
+configs = list(itertools.product(SIZE_LIST))
+
+
+@triton.testing.perf_report(
+    triton.testing.Benchmark(
+        x_names=["size"],
+        x_vals=configs,
+        line_arg="provider",
+        line_vals=["jit", "torch"],
+        line_names=["SGL JIT Kernel", "PyTorch"],
+        styles=[("blue", "-"), ("red", "--")],
+        ylabel="us",
+        plot_name="scale-performance",
+        args={},
+    )
+)
+def benchmark(size: int, provider: str):
+    src = torch.randn(size, dtype=DEFAULT_DTYPE, device=DEFAULT_DEVICE)
+    factor = 2.0
+
+    if provider == "jit":
+        fn = lambda: jit_scale(src, factor)
+    else:
+        fn = lambda: src * factor
+
+    return run_benchmark(fn)
+
+
+if __name__ == "__main__":
+    benchmark.run(print_data=True)
+```
+
+Run locally:
+
+```bash
+python python/sglang/jit_kernel/benchmark/bench_scale.py
+```
+
+Run the benchmark suite the way CI does:
+
+```bash
+cd test && python3 run_suite.py --hw cuda --suite stage-b-kernel-benchmark-1-gpu-large
+```
+
+---
+
+## Troubleshooting
+
+- **`No CI registry found in ...` from `run_suite.py`**: add a module-level `register_cuda_ci(...)` with literal `est_time` and `suite` (and optional `nightly=True`); starred args and non-literal values break AST collection
+- **JIT compilation fails**: ensure the `.cuh` file is under `python/sglang/jit_kernel/csrc/`; reduce template argument combinations
+- **CUDA crash / illegal memory access**: `CUDA_LAUNCH_BLOCKING=1`; `compute-sanitizer --tool memcheck python ...`
+- **Unstable benchmark results**: `run_benchmark` uses CUDA-graph-based timing by default
+
+---
+
+## References
+
+- `docs/developer_guide/development_jit_kernel_guide.md`
+- `test/run_suite.py` — suite names, discovery of `jit_kernel/tests/` and `jit_kernel/benchmark/`, execution entrypoint for CI
+- `python/sglang/test/ci/ci_register.py` — `register_cuda_ci` and AST registration rules
+- `python/sglang/jit_kernel/utils.py` — `cache_once`, `load_jit`, `make_cpp_args`, `should_run_full_tests`, `get_ci_test_range`
+- `python/sglang/jit_kernel/include/sgl_kernel/tensor.h` — `TensorMatcher`, `SymbolicSize/DType/Device`
+- `python/sglang/jit_kernel/include/sgl_kernel/utils.cuh` — type aliases, `LaunchKernel`, `SGL_DEVICE`
+- `python/sglang/jit_kernel/include/sgl_kernel/vec.cuh` — `AlignedVector`
+- `python/sglang/jit_kernel/include/sgl_kernel/tile.cuh` — `tile::Memory`
+- `python/sglang/jit_kernel/include/sgl_kernel/type.cuh` — `dtype_trait`, `packed_t`, `device::cast`
+- `python/sglang/jit_kernel/include/sgl_kernel/math.cuh` — `device::math::`
+- `python/sglang/jit_kernel/include/sgl_kernel/warp.cuh` — `warp::reduce_sum/max`
+- `python/sglang/jit_kernel/include/sgl_kernel/cta.cuh` — `cta::reduce_max`
+- `python/sglang/jit_kernel/include/sgl_kernel/atomic.cuh` — `atomic::max`
+- `python/sglang/jit_kernel/include/sgl_kernel/runtime.cuh` — occupancy / SM count helpers
+- `python/sglang/jit_kernel/csrc/add_constant.cuh` — minimal runnable reference
+- `python/sglang/jit_kernel/csrc/elementwise/rmsnorm.cuh` — real example using `TensorMatcher` + `LaunchKernel` + `tile::Memory`
+- `python/sglang/jit_kernel/csrc/elementwise/qknorm.cuh` — real example using `runtime::get_blocks_per_sm` + persistent kernel pattern
+- `python/sglang/jit_kernel/benchmark/utils.py` — benchmark helpers
+
+## Summary of Files Created
+
+```
+python/sglang/jit_kernel/csrc/elementwise/scale.cuh   # NEW: CUDA kernel
+python/sglang/jit_kernel/scale.py                     # NEW: Python wrapper
+python/sglang/jit_kernel/tests/test_scale.py          # NEW: Tests
+python/sglang/jit_kernel/benchmark/bench_scale.py     # NEW: Benchmark
+```

third_party/sglang/.claude/skills/add-sgl-kernel/SKILL.md

@@ -0,0 +1,363 @@
+---
+name: add-sgl-kernel
+description: Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)
+---
+
+# Tutorial: Adding a New Kernel to `sgl-kernel` (AOT / Heavyweight)
+
+This tutorial walks through adding a simple element-wise scale operation as an AOT kernel. We'll implement `scale(x, factor) = x * factor` to demonstrate the complete workflow.
+
+## Goal
+
+Add a new operation that scales each element of a tensor by a scalar factor:
+
+- Input: tensor `x` (CUDA) and scalar `factor` (float)
+- Output: `x * factor` (element-wise, in-place or into pre-allocated `out`)
+- Supported dtypes: **FP16 (`torch.float16`), BF16 (`torch.bfloat16`), FP32 (`torch.float32`)**
+  - Dispatched via `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro (defined in `sgl-kernel/include/utils.h`)
+
+## Two rules of thumb (must follow)
+
+1. **Prefer `python/sglang/jit_kernel` first** when the kernel does **not** depend on CUTLASS or another large C++ project. This is the default path for lightweight kernels that benefit from rapid iteration.
+2. **Prefer `sgl-kernel`** when the kernel **does** depend on CUTLASS or another large C++ project, or when it should be part of the AOT wheel / torch op registration flow.
+3. **Exception**: if the dependency is `flashinfer`, or CUTLASS that is already provided through `flashinfer`, the kernel can still be implemented as `jit_kernel`.
+
+In addition, every new kernel must ship with:
+
+- **Tests** (pytest)
+- **A benchmark script** (triton.testing)
+
+---
+
+## Repository integration map
+
+You will typically touch these files/areas:
+
+- Implementation: `sgl-kernel/csrc/elementwise/scale.cu` (pick the right subdirectory)
+- Public declarations: `sgl-kernel/include/sgl_kernel_ops.h`
+- Torch extension registration: `sgl-kernel/csrc/common_extension.cc`
+- Build: `sgl-kernel/CMakeLists.txt` (`set(SOURCES ...)`)
+- Python API: `sgl-kernel/python/sgl_kernel/` and `sgl-kernel/python/sgl_kernel/__init__.py`
+- Tests: `sgl-kernel/tests/test_scale.py`
+- Benchmarks: `sgl-kernel/benchmark/bench_scale.py`
+
+---
+
+## Step 1: Implement the kernel in `csrc/`
+
+Pick the right subdirectory:
+
+- `csrc/elementwise/` — for element-wise ops (our example)
+- `csrc/gemm/`, `csrc/attention/`, `csrc/moe/` — for other categories
+
+Create `sgl-kernel/csrc/elementwise/scale.cu`:
+
+```cpp
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <torch/all.h>
+
+#include "utils.h"  // DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16
+
+// scale_kernel: out[i] = input[i] * factor
+// Supports float, half (__half), __nv_bfloat16 via template T
+template <typename T>
+__global__ void scale_kernel(T* __restrict__ out,
+                              const T* __restrict__ input,
+                              float factor,
+                              int64_t n) {
+  int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+  if (idx < n) {
+    out[idx] = static_cast<T>(static_cast<float>(input[idx]) * factor);
+  }
+}
+
+void scale(at::Tensor& out, const at::Tensor& input, double factor) {
+  TORCH_CHECK(input.is_cuda(),       "input must be a CUDA tensor");
+  TORCH_CHECK(input.is_contiguous(), "input must be contiguous");
+  TORCH_CHECK(out.is_cuda(),         "out must be a CUDA tensor");
+  TORCH_CHECK(out.is_contiguous(),   "out must be contiguous");
+  TORCH_CHECK(out.sizes() == input.sizes(),  "out and input must have the same shape");
+  TORCH_CHECK(out.scalar_type() == input.scalar_type(),
+              "out and input must have the same dtype");
+
+  const int64_t n = input.numel();
+  const int threads = 256;
+  const int blocks  = (n + threads - 1) / threads;
+
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
+
+  // Dispatches over float, float16, bfloat16
+  DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16(input.scalar_type(), c_type, [&] {
+    scale_kernel<c_type><<<blocks, threads, 0, stream>>>(
+        static_cast<c_type*>(out.data_ptr()),
+        static_cast<const c_type*>(input.data_ptr()),
+        static_cast<float>(factor),
+        n);
+    cudaError_t status = cudaGetLastError();
+    TORCH_CHECK(status == cudaSuccess,
+                "scale_kernel launch failed: ", cudaGetErrorString(status));
+    return true;
+  });
+}
+```
+
+**Key points:**
+
+- Use `at::Tensor` (PyTorch tensors), `TORCH_CHECK` for validation, `at::cuda::getCurrentCUDAStream()` for stream
+- Keep Python wrappers thin; do shape/dtype/device validation in C++ right around the launch path
+- `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` covers `float`, `half` (FP16), `__nv_bfloat16` (BF16)
+- Add device error checking after every kernel launch
+- If a kernel only works on certain architectures, enforce that with `TORCH_CHECK` and skip logic in tests
+
+---
+
+## Step 2: Add a C++ declaration in `include/sgl_kernel_ops.h`
+
+Edit `sgl-kernel/include/sgl_kernel_ops.h`, add to the elementwise section:
+
+```cpp
+void scale(at::Tensor& out, const at::Tensor& input, double factor);
+```
+
+---
+
+## Step 3: Register the op in `csrc/common_extension.cc`
+
+Edit `sgl-kernel/csrc/common_extension.cc`, inside `TORCH_LIBRARY_FRAGMENT(sgl_kernel, m)`:
+
+```cpp
+// From csrc/elementwise
+m.def("scale(Tensor! out, Tensor input, float factor) -> ()");
+m.impl("scale", torch::kCUDA, &scale);
+```
+
+**Key points:**
+
+- `Tensor!` means in-place / mutable output argument
+- The schema is important for `torch.compile` and for consistent call signatures
+- Keep the torch schema in PyTorch scalar types (`float` here), but note that the C++ launcher signature still needs `double` for scalar arguments accepted by `torch::Library`
+
+---
+
+## Step 4: Add the new source file to `CMakeLists.txt`
+
+Edit `sgl-kernel/CMakeLists.txt`, add to `set(SOURCES ...)`:
+
+```cmake
+csrc/elementwise/scale.cu
+```
+
+**Key points:**
+
+- Keep the list **alphabetically sorted** (the file explicitly requires this)
+- If the kernel has arch constraints, reflect that in tests/benchmarks via skip logic
+
+---
+
+## Step 5: Expose a Python API under `sgl-kernel/python/sgl_kernel/`
+
+Prefer following the existing module organization first. For elementwise kernels, the usual pattern is:
+
+- implement the Python wrapper in `sgl-kernel/python/sgl_kernel/elementwise.py`
+- then re-export it from `sgl-kernel/python/sgl_kernel/__init__.py`
+
+For example, in `sgl-kernel/python/sgl_kernel/elementwise.py`, add:
+
+```python
+import torch
+
+def scale(
+    input: torch.Tensor,
+    factor: float,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """
+    Element-wise scale: out = input * factor.
+
+    Supported dtypes: torch.float16, torch.bfloat16, torch.float32.
+
+    Parameters
+    ----------
+    input  : CUDA input tensor
+    factor : scale factor (float)
+    out    : optional pre-allocated CUDA output tensor (same shape/dtype as input)
+    """
+    if out is None:
+        out = torch.empty_like(input)
+    torch.ops.sgl_kernel.scale.default(out, input, factor)
+    return out
+```
+
+Then re-export it from `sgl-kernel/python/sgl_kernel/__init__.py` following the existing import style used by other kernels.
+
+---
+
+## Step 6: Write tests (required)
+
+Create `sgl-kernel/tests/test_scale.py`:
+```python
+import pytest
+
+import torch
+import sgl_kernel
+
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float32])
+@pytest.mark.parametrize("size", [128, 1024, 4096, 65536])
+@pytest.mark.parametrize("factor", [0.5, 1.0, 2.0])
+def test_scale_correctness(dtype, size, factor):
+    input = torch.randn(size, dtype=dtype, device="cuda")
+    out   = torch.empty_like(input)
+
+    result = sgl_kernel.scale(input, factor, out=out)
+    assert result is out
+
+    expected = input * factor
+    rtol, atol = (1e-5, 1e-6) if dtype == torch.float32 else (1e-2, 1e-2)
+    torch.testing.assert_close(out, expected, rtol=rtol, atol=atol)
+
+
+def test_scale_shape_mismatch():
+    input = torch.randn(128, dtype=torch.float16, device="cuda")
+    out   = torch.empty(256, dtype=torch.float16, device="cuda")
+    with pytest.raises(RuntimeError, match="same shape"):
+        sgl_kernel.scale(input, 2.0, out=out)
+
+
+def test_scale_cpu_input():
+    input = torch.randn(128, dtype=torch.float16)  # CPU
+    out   = torch.empty_like(input)
+    with pytest.raises(RuntimeError, match="CUDA"):
+        sgl_kernel.scale(input, 2.0, out=out)
+
+
+if __name__ == "__main__":
+    import sys
+    sys.exit(pytest.main([__file__, "-q"]))
+```
+
+---
+
+## Step 7: Add a benchmark (required)
+
+Create `sgl-kernel/benchmark/bench_scale.py`:
+
+```python
+import itertools
+
+import torch
+import triton
+import triton.testing
+
+import sgl_kernel
+from sglang.utils import is_in_ci
+
+IS_CI = is_in_ci()
+
+dtypes  = [torch.float16] if IS_CI else [torch.float16, torch.bfloat16, torch.float32]
+sizes   = [4096] if IS_CI else [2**n for n in range(10, 20)]  # 1K … 512K
+factors = [2.0]
+
+configs = list(itertools.product(dtypes, sizes))
+
+
+def torch_scale(input: torch.Tensor, factor: float) -> torch.Tensor:
+    return input * factor
+
+
+@triton.testing.perf_report(
+    triton.testing.Benchmark(
+        x_names=["dtype", "size"],
+        x_vals=configs,
+        line_arg="provider",
+        line_vals=["sglang", "torch"],
+        line_names=["SGL Kernel", "PyTorch"],
+        styles=[("green", "-"), ("red", "--")],
+        ylabel="µs (median)",
+        plot_name="scale-performance",
+        args={},
+    )
+)
+def benchmark(dtype, size, provider):
+    input  = torch.randn(size, dtype=dtype, device="cuda")
+    out    = torch.empty_like(input)
+    factor = 2.0
+
+    if provider == "sglang":
+        fn = lambda: sgl_kernel.scale(input, factor, out=out)
+    else:
+        fn = lambda: torch_scale(input, factor)
+
+    ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+        fn, quantiles=[0.5, 0.2, 0.8]
+    )
+    return 1000 * ms, 1000 * max_ms, 1000 * min_ms
+
+
+if __name__ == "__main__":
+    benchmark.run(print_data=True)
+```
+
+---
+
+## Step 8: Build
+
+Build:
+
+```bash
+cd sgl-kernel
+make build -j16
+```
+
+If you need to limit host resource usage:
+
+```bash
+cd sgl-kernel
+make build -j1 MAX_JOBS=2 CMAKE_ARGS="-DSGL_KERNEL_COMPILE_THREADS=1"
+```
+
+---
+
+## Step 9: Validate
+
+After building successfully, run the test and benchmark:
+
+```bash
+pytest sgl-kernel/tests/test_scale.py -q
+python sgl-kernel/benchmark/bench_scale.py
+```
+
+---
+
+## Troubleshooting
+
+- **Async CUDA errors**: `CUDA_LAUNCH_BLOCKING=1`
+- **Memory errors**: `compute-sanitizer --tool memcheck python ...`
+- **Build is too slow / OOM**: reduce `MAX_JOBS` and `SGL_KERNEL_COMPILE_THREADS`
+- **Binary bloat**: use `sgl-kernel/analyze_whl_kernel_sizes.py`
+- **CMake sources list**: if your `.cu` file is missing from `SOURCES`, the symbol will be undefined at link time
+
+---
+
+## References
+
+- `sgl-kernel/README.md`
+- `sgl-kernel/include/sgl_kernel_ops.h`
+- `sgl-kernel/csrc/common_extension.cc`
+- `sgl-kernel/CMakeLists.txt`
+- `sgl-kernel/include/utils.h` — `DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16` macro and friends
+- `sgl-kernel/csrc/elementwise/activation.cu` — reference for the FP16/BF16/FP32 dispatch pattern
+
+## Summary of Files Created/Modified
+
+```
+sgl-kernel/csrc/elementwise/scale.cu          # NEW: CUDA kernel + launcher
+sgl-kernel/include/sgl_kernel_ops.h           # MODIFIED: C++ declaration
+sgl-kernel/csrc/common_extension.cc           # MODIFIED: schema + dispatch registration
+sgl-kernel/CMakeLists.txt                     # MODIFIED: add source file (alphabetical)
+sgl-kernel/python/sgl_kernel/elementwise.py   # MODIFIED: Python wrapper
+sgl-kernel/python/sgl_kernel/__init__.py      # MODIFIED: re-export Python API
+sgl-kernel/tests/test_scale.py                # NEW: tests
+sgl-kernel/benchmark/bench_scale.py           # NEW: benchmark
+```

third_party/sglang/.claude/skills/ci-workflow-guide/SKILL.md

@@ -0,0 +1,386 @@
+---
+name: ci-workflow-guide
+description: Guide to SGLang CI workflow orchestration — stage ordering, fast-fail, gating, partitioning, execution modes, and debugging CI failures. Use when modifying CI workflows, adding stages, debugging CI pipeline issues, or understanding how tests are dispatched and gated across stages.
+---
+
+# SGLang CI Workflow Orchestration Guide
+
+This skill covers the CI **infrastructure** layer — how tests are dispatched, gated, and fast-failed across stages. For test authoring (templates, fixtures, registration, model selection), see the [write-sglang-test skill](../write-sglang-test/SKILL.md).
+
+---
+
+## Naming Conventions
+
+- **Suite**: `stage-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `stage-b-test-1-gpu-small`)
+- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`)
+
+---
+
+## Key Files
+
+| File | Role |
+|------|------|
+| `.github/workflows/pr-test.yml` | Main workflow — all stages, jobs, conditions, matrix definitions |
+| `.github/workflows/pr-gate.yml` | PR gating: draft check, `run-ci` label, per-user rate limiting |
+| `.github/actions/check-stage-health/action.yml` | Cross-job fast-fail: queries API for any failed job |
+| `.github/actions/wait-for-jobs/action.yml` | Stage gating: polls API until stage jobs complete |
+| `.github/actions/check-maintenance/action.yml` | Maintenance mode check |
+| `test/run_suite.py` | Suite runner: collects, filters, partitions, executes tests |
+| `python/sglang/test/ci/ci_register.py` | Test registration (AST-parsed markers), LPT auto-partition |
+| `python/sglang/test/ci/ci_utils.py` | `run_unittest_files()`: execution, retry, continue-on-error |
+| `scripts/ci/utils/slash_command_handler.py` | Handles slash commands from PR comments |
+
+---
+
+## Architecture Overview
+
+```
+ ┌──────────────┐
+ │ build kernel │
+ └──────┬───────┘
+        │
+        ├─ check-changes ──── detects which packages changed
+        │                      (main_package, sgl_kernel, jit_kernel, multimodal_gen)
+        │
+        ├─ call-gate ──────── pr-gate.yml (draft? label? rate limit?)
+        │
+        ├─────────────────────────────────────────────────────┐
+        │                                                     │
+        ▼                                                     │
+ ┌─────────────────────────────────────┐                      │
+ │          Stage A (~3 min)           │                      │
+ │         pre-flight check            │                      │
+ │                                     │                      │
+ │  ┌─────────────────────────────┐    │                      │
+ │  │ stage-a-test-1-gpu-small    │    │                      │
+ │  │ (small GPUs)                │    │                      │
+ │  └─────────────────────────────┘    │                      │
+ │  ┌─────────────────────────────┐    │                      │
+ │  │ stage-a-test-cpu            │    │                      │
+ │  │ (CPU)                       │    │                      │
+ │  └─────────────────────────────┘    │                      │
+ └──────┬──────────────────────────────┘                      │
+        │                                                     │
+        ▼                                                     ▼
+ ┌─────────────────────────────────────┐          ┌──────────────────────────┐
+ │          Stage B (~30 min)          │          │      kernel test         │
+ │           basic tests               │          └──────────────────────────┘
+ │                                     │          ┌──────────────────────────┐
+ │  ┌─────────────────────────────┐    │          │   multimodal gen test    │
+ │  │ stage-b-test-1-gpu-small    │    │          └──────────────────────────┘
+ │  │ (small GPUs, e.g. 5090)     │    │
+ │  └─────────────────────────────┘    │
+ │  ┌─────────────────────────────┐    │
+ │  │ stage-b-test-1-gpu-large    │    │
+ │  │ (large GPUs, e.g. H100)     │    │
+ │  └─────────────────────────────┘    │
+ │  ┌─────────────────────────────┐    │
+ │  │ stage-b-test-2-gpu-large    │    │
+ │  │ (large GPUs, e.g. H100)     │    │
+ │  └─────────────────────────────┘    │
+ └──────┬──────────────────────────────┘
+        │
+        ▼
+ ┌─────────────────────────────────────┐
+ │          Stage C (~30 min)          │
+ │          advanced tests             │
+ │                                     │
+ │  ┌─────────────────────────────┐    │
+ │  │ stage-c-test-4-gpu-h100     │    │
+ │  │ (H100 GPUs)                 │    │
+ │  └─────────────────────────────┘    │
+ │  ┌─────────────────────────────┐    │
+ │  │ stage-c-test-8-gpu-h200     │    │
+ │  │ (8 x H200 GPUs)             │    │
+ │  └─────────────────────────────┘    │
+ │  ┌─────────────────────────────┐    │
+ │  │ stage-c-test-4-gpu-b200     │    │
+ │  │ (4 x B200 GPUs)             │    │
+ │  └─────────────────────────────┘    │
+ │  ┌─────────────────────────────┐    │
+ │  │ Other advanced tests        │    │
+ │  │ (DeepEP, PD Disagg, GB300)  │    │
+ │  └─────────────────────────────┘    │
+ └──────┬──────────────────────────────┘
+        │
+        ▼
+ ┌─────────────────────────────────────┐
+ │         pr-test-finish              │
+ │  aggregates all results, fails if   │
+ │  any job failed/cancelled           │
+ └─────────────────────────────────────┘
+```
+
+**Every stage test job** includes a `check-stage-health` step after checkout — if any job in the run has already failed, the job fast-fails (red X) with a root cause annotation.
+
+**Scheduled runs** skip `wait-for-stage-*` jobs, running all stages in parallel. Fast-fail is also disabled.
+
+---
+
+## Fast-Fail Layers
+
+4 layers of fast-fail, from fine to coarse:
+
+| Layer | Mechanism | Granularity | Disabled on schedule? |
+|-------|-----------|-------------|----------------------|
+| **1. Test method → file** | `unittest -f` (failfast) | One test method fails → entire test file stops immediately | Yes |
+| **2. File → suite** | `run_unittest_files()` default | One test file fails → entire suite stops (`--continue-on-error` off) | Yes |
+| **3. Job → job (same stage)** | `check-stage-health` action | One job fails → other waiting jobs in same stage fast-fail (red X) | Yes |
+| **4. Stage → stage (cross-stage)** | `wait-for-stage` + `needs` | Stage A fails → stage B/C jobs skip entirely (never get a runner) | Yes (wait jobs skipped) |
+
+- **Layer 1**: `-f` flag appended to all `python3 -m pytest` / `unittest` invocations in `ci_utils.py`
+- **Layer 2**: `--continue-on-error` flag in `run_suite.py` — off for PRs, on for scheduled runs
+- **Layer 3**: `check-stage-health` auto-detects `schedule` event and skips; filters out cascade failures to show only root cause jobs
+- **Layer 4**: `wait-for-stage-*` jobs are conditioned on `github.event_name == 'pull_request'` — skipped for scheduled runs
+
+---
+
+## Execution Modes
+
+| Aspect | PR (`pull_request`) | Scheduled (`cron`, every 6h) | `/rerun-stage` (`workflow_dispatch`) |
+|--------|---------------------|------------------------------|--------------------------------------|
+| **Stage ordering** | Sequential: A → B → C via `wait-for-stage-*` | Parallel (all at once) | Single target stage only |
+| **Cross-job fast-fail** | Yes (`check-stage-health`) | Yes | Yes |
+| **continue-on-error** | No (stop at first failure within suite) | Yes (run all tests) | No |
+| **Retry** | Enabled | Enabled | Enabled |
+| **max_parallel** | 3 (default), 14 if `high priority` label | 14 | 3 (default), 14 if `high priority` |
+| **PR gate** | Yes (draft, label, rate limit) | Skipped | Skipped |
+| **Concurrency** | `cancel-in-progress: true` per branch | Queue (no cancel) | Isolated per stage+SHA |
+
+---
+
+## Stage Gating (`wait-for-jobs` action)
+
+`wait-for-stage-a` and `wait-for-stage-b` are lightweight `ubuntu-latest` jobs that poll the GitHub Actions API.
+
+**How it works:**
+1. Calls `listJobsForWorkflowRun` to list all jobs in the current run
+2. Matches jobs by exact name or prefix (for matrix jobs, e.g., `stage-b-test-1-gpu-small (3)`)
+3. If any matched job has `conclusion === 'failure'` → fail immediately (fast-fail)
+4. If all matched jobs are completed and count matches `expected_count` → success
+5. Otherwise → sleep `poll-interval-seconds` (default: 60s) and retry
+6. Timeout after `max-wait-minutes` (240 min for stage-a, 480 min for stage-b)
+
+**Job specs example** (stage-b):
+```json
+[
+  {"prefix": "stage-b-test-1-gpu-small", "expected_count": 8},
+  {"prefix": "stage-b-test-1-gpu-large", "expected_count": 14},
+  {"prefix": "stage-b-test-2-gpu-large", "expected_count": 4},
+  {"prefix": "stage-b-test-4-gpu-b200", "expected_count": 1}
+]
+```
+
+> **Critical**: `expected_count` must match the matrix size. If you add/remove matrix entries, update the wait job's spec accordingly.
+
+**PR only**: Condition `github.event_name == 'pull_request' && !inputs.target_stage` — scheduled runs and `/rerun-stage` skip these entirely, allowing parallel execution.
+
+---
+
+## Cross-Job Fast-Fail (`check-stage-health` action)
+
+Composite action called after checkout in every stage test job (21 jobs total across `pr-test.yml`, `pr-test-multimodal-gen.yml`, `pr-test-sgl-kernel.yml`, `pr-test-jit-kernel.yml`).
+
+**How it works:**
+1. Queries `listJobsForWorkflowRun` for the current workflow run
+2. Filters for **root cause failures only** — jobs with `conclusion === 'failure'` whose failing step is NOT `check-stage-health` (excludes cascade failures)
+3. If root cause failures found → calls `core.setFailed()` with the list of root cause job names
+4. If none → does nothing (step succeeds)
+
+**Cascade filtering**: When job A fast-fails due to health check, it also has `conclusion: failure`. Without filtering, job B would list both the original failure AND job A's fast-fail. The filter checks each failed job's `steps` array — if the failing step name contains `check-stage-health` or `Check stage health`, it's excluded from the root cause list.
+
+**Usage pattern:**
+```yaml
+steps:
+  - name: Checkout code
+    uses: actions/checkout@v4
+    ...
+
+  - uses: ./.github/actions/check-stage-health
+    id: stage-health
+
+  - name: Install dependencies        # skipped automatically if health check failed
+    ...                                # (default if: success() is false)
+
+  - name: Run test                     # also skipped
+    ...
+```
+
+**Visual effect**: Job shows **red X** (failure) with error annotation showing root cause job names. Subsequent steps are naturally skipped (default `if: success()` is false after a failed step). No per-step `if` guards needed.
+
+**No stage filtering**: Checks ALL jobs in the run, not just the current stage. Any failure anywhere triggers fast-fail.
+
+**Error message example:**
+```
+Fast-fail: skipping — root cause job(s): stage-b-test-1-gpu-small (0), stage-b-test-1-gpu-small (1)
+```
+
+---
+
+## Within-Suite Failure Handling
+
+Controlled by `run_unittest_files()` in `python/sglang/test/ci/ci_utils.py`.
+
+### Flags
+
+| Flag | PR default | Scheduled default | Effect |
+|------|------------|-------------------|--------|
+| `--continue-on-error` | Off | On | Off: stop at first failure. On: run all files, report all failures at end |
+| `--enable-retry` | On | On | Retry retriable failures (accuracy/perf assertions) |
+| `--max-attempts` | 2 | 2 | Max attempts per file including initial run |
+
+### Retry Classification
+
+When a test fails and retry is enabled, the output is classified:
+
+**Non-retriable** (checked first — real code errors):
+`SyntaxError`, `ImportError`, `ModuleNotFoundError`, `NameError`, `TypeError`, `AttributeError`, `RuntimeError`, `CUDA out of memory`, `OOM`, `Segmentation fault`, `core dumped`, `ConnectionRefusedError`, `FileNotFoundError`
+
+**Retriable** (accuracy/performance):
+`AssertionError` with comparison patterns (`not greater than`, `not less than`, `not equal to`), `accuracy`, `score`, `latency`, `throughput`, `timeout`
+
+**Default**: Unknown `AssertionError` → retriable. Other unknown failures → not retriable.
+
+### How `continue_on_error` is set
+
+In `pr-test.yml`'s `check-changes` job:
+- `schedule` runs or `run_all_tests` flag → `continue_on_error = 'true'`
+- PR runs → `continue_on_error = 'false'`
+
+Each test job propagates via:
+```yaml
+env:
+  CONTINUE_ON_ERROR_FLAG: ${{ needs.check-changes.outputs.continue_on_error == 'true' && '--continue-on-error' || '' }}
+run: |
+  python3 run_suite.py --hw cuda --suite <name> $CONTINUE_ON_ERROR_FLAG
+```
+
+---
+
+## Test Partitioning
+
+Large suites are split across matrix jobs using the **LPT (Longest Processing Time) heuristic** in `ci_register.py:auto_partition()`:
+
+1. Sort tests by `est_time` descending, filename as tie-breaker (deterministic)
+2. Greedily assign each test to the partition with smallest cumulative time
+3. Result: roughly equal total time per partition
+
+**Partition table** (CUDA per-commit suites):
+
+| Suite | Partitions | Runner | max_parallel |
+|-------|-----------|--------|-------------|
+| `stage-a-test-1-gpu-small` | 1 (no matrix) | `1-gpu-5090` | — |
+| `stage-a-test-cpu` | 1 (no matrix) | `ubuntu-latest` | — |
+| `stage-b-test-1-gpu-small` | 8 | `1-gpu-5090` | 8 |
+| `stage-b-test-1-gpu-large` | 14 | `1-gpu-h100` | dynamic (3 or 14) |
+| `stage-b-test-2-gpu-large` | 4 | `2-gpu-h100` | — |
+| `stage-b-test-4-gpu-b200` | 1 (no matrix) | `4-gpu-b200` | — |
+| `stage-b-kernel-unit-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |
+| `stage-b-kernel-unit-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — |
+| `stage-b-kernel-benchmark-1-gpu-large` | 1 (no matrix) | `1-gpu-h100` | — |
+| `stage-c-test-4-gpu-h100` | 3 | `4-gpu-h100` | — |
+| `stage-c-test-8-gpu-h200` | 4 | `8-gpu-h200` | — |
+| `stage-c-test-8-gpu-h20` | 2 | `8-gpu-h20` | — |
+| `stage-c-test-deepep-4-gpu-h100` | 1 (no matrix) | `4-gpu-h100` | — |
+| `stage-c-test-deepep-8-gpu-h200` | 1 (no matrix) | `8-gpu-h200` | — |
+| `stage-c-test-4-gpu-b200` | 4 | `4-gpu-b200` | — |
+| `stage-c-test-4-gpu-gb200` | 1 (no matrix) | `4-gpu-gb200` | — |
+
+> **Note**: Kernel suites (`stage-b-kernel-*`) run via `pr-test-jit-kernel.yml` and `pr-test-sgl-kernel.yml`, not the main `pr-test.yml`. Multimodal diffusion uses `python/sglang/multimodal_gen/test/run_suite.py`, not `test/run_suite.py`.
+
+**Workflow usage:**
+```yaml
+strategy:
+  matrix:
+    partition: [0, 1, 2, 3, 4, 5, 6, 7]
+steps:
+  - run: python3 run_suite.py --hw cuda --suite stage-b-test-1-gpu-small \
+           --auto-partition-id ${{ matrix.partition }} --auto-partition-size 8
+```
+
+---
+
+## check-changes Job
+
+Determines which test suites to run based on file changes.
+
+### Detection Methods
+
+| Trigger | Method | Details |
+|---------|--------|---------|
+| `pull_request` | `dorny/paths-filter` | Detects changes via GitHub diff |
+| `workflow_dispatch` (with `pr_head_sha`) | GitHub API | `repos/{repo}/compare/main...{sha}` |
+| `schedule` / `run_all_tests` | Force all true | Runs everything |
+
+### Output Flags
+
+| Output | Triggers |
+|--------|----------|
+| `main_package` | Stage A/B/C test suites |
+| `sgl_kernel` | Kernel wheel builds + kernel test suites |
+| `jit_kernel` | JIT kernel test workflow |
+| `multimodal_gen` | Multimodal-gen test workflow |
+
+> **Note**: `sgl_kernel` is forced to `false` when `target_stage` is set, because `sgl-kernel-build-wheels` won't run and wheel artifacts won't be available.
+
+---
+
+## Concurrency Control
+
+```
+group: pr-test-{event_name}-{branch}-{pr_sha}-{stage}
+```
+
+| Segment | Source | Purpose |
+|---------|--------|---------|
+| `event_name` | `github.event_name` | Prevents scheduled runs colliding with fork PRs named `main` |
+| `branch` | `github.head_ref \|\| github.ref_name` | Per-branch isolation |
+| `pr_sha` | `inputs.pr_head_sha \|\| 'current'` | Isolates `/rerun-stage` from main runs |
+| `stage` | `inputs.target_stage \|\| 'all'` | Allows parallel stage dispatches |
+
+`cancel-in-progress: true` for `pull_request` events (new push cancels old run), `false` for `workflow_call`.
+
+---
+
+## How To: Add a New Stage Job
+
+1. Define the job in `pr-test.yml` with `needs: [check-changes, call-gate, wait-for-stage-X, ...]`
+2. Copy the `if:` condition pattern from an existing same-stage job (handles `target_stage`, `schedule`, `main_package`)
+3. Add `checkout` step
+4. Add `check-stage-health` step (after checkout) — if any prior job failed, `core.setFailed()` fires and all subsequent steps auto-skip via default `if: success()`
+5. Add `check-maintenance` step
+6. Add `download-artifact` step if `sgl_kernel` changed
+7. Add `install dependencies` step
+8. Add `run test` step with `$CONTINUE_ON_ERROR_FLAG`
+9. Add `upload-cuda-coredumps` step with `if: always()`
+10. Register the suite name in `PER_COMMIT_SUITES` in `test/run_suite.py`
+11. If using matrix, add `--auto-partition-id` and `--auto-partition-size` to the run command
+12. **Update `wait-for-stage-X`** job spec with the new job name and `expected_count` (if matrix)
+13. **Add the job to `pr-test-finish.needs`** list
+
+---
+
+## How To: Debug CI Failures
+
+| Symptom | Likely cause | What to check |
+|---------|-------------|---------------|
+| All stage-B/C jobs green but steps skipped | Earlier job failed, `check-stage-health` triggered | Find the actual failed job (red X) |
+| `wait-for-stage-b` timeout | `expected_count` doesn't match matrix size | Verify job spec counts match `matrix:` array length |
+| `pr-test-finish` fails but all jobs green | A job was `cancelled` (counts as failure in finish) | Check concurrency cancellation |
+| Tests pass locally but fail in CI | Partition assignment, runner GPU type, or `est_time` inaccuracy | Check which partition the test lands in; verify runner label |
+| Flaky test retried and passed | Retriable failure (accuracy/perf) | Check `[CI Retry]` markers in job logs |
+| Flaky test NOT retried | Matched non-retriable pattern | Check if error matches `NON_RETRIABLE_PATTERNS` in `ci_utils.py` |
+
+---
+
+## Slash Commands
+
+| Command | Effect |
+|---------|--------|
+| `/tag-run-ci-label` | Adds `run-ci` label to PR |
+| `/rerun-failed-ci` | Reruns failed jobs in the latest workflow run |
+| `/tag-and-rerun-ci` | Adds label + reruns |
+| `/rerun-stage <stage>` | Dispatches `pr-test.yml` with `target_stage=<stage>` |
+| `/rerun-test <test-file>` | Reruns a specific test file via `rerun-test.yml` |
+
+Handled by `scripts/ci/utils/slash_command_handler.py` → `.github/workflows/slash-command-handler.yml`.

third_party/sglang/.claude/skills/debug-cuda-crash/SKILL.md

@@ -0,0 +1,657 @@
+---
+name: debug-cuda-crash
+description: Call this skill when you need to debug CUDA crashes in SGLang using kernel API logging
+---
+
+# Tutorial: Debugging CUDA Crashes with Kernel API Logging
+
+This tutorial shows you how to debug CUDA crashes and errors in SGLang using the `@debug_kernel_api` logging decorator.
+
+## Goal
+
+When your code crashes with CUDA errors such as illegal memory access, device-side assert, out-of-bounds, or NaN/Inf, use kernel API logging to:
+- Capture input tensors BEFORE the crash occurs
+- Understand what data caused the problem
+- Track tensor shapes, dtypes, and values through the call boundary that triggered the crash
+- Detect numerical issues such as NaN, Inf, or obviously wrong shapes
+
+## Why Use Kernel API Logging?
+
+**Problem**: CUDA errors often crash the program before normal debugging output is flushed.
+
+**Solution**: SGLang's `@debug_kernel_api` decorator logs inputs before execution, so you can still see what caused the crash even after the program aborts.
+
+## What Is Covered?
+
+The current logging coverage focuses on the highest-value kernel boundaries in SGLang:
+- Custom ops registered through `register_custom_op(...)`
+- External custom ops registered through `register_custom_op_from_extern(...)`
+- LLM attention, linear, quantization, and multi-platform wrapper entry points
+- Diffusion attention impl, linear, rotary, and custom-op wrapper entry points
+- Selected direct `torch.ops.sglang.*` hotspots and model-specific bypasses
+
+This means the logging is useful for both LLM and diffusion kernel debugging, but it does not automatically cover every pure PyTorch call in the repository.
+
+## Step 1: Enable Kernel API Logging
+
+### Basic Logging (Function Names Only)
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=1
+export SGLANG_KERNEL_API_LOGDEST=stdout
+
+python my_script.py
+```
+
+Output:
+```
+================================================================================
+[2026-03-19 00:47:06] SGLang Kernel API Call: RMSNorm.forward
+================================================================================
+[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply
+================================================================================
+[2026-03-19 00:47:06] SGLang Kernel API Call: sglang.custom_op.fused_inplace_qknorm
+```
+
+This is a real level-1 excerpt captured from `Qwen/Qwen3-0.6B`.
+
+### Detailed Logging (Inputs with Metadata)
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+export SGLANG_KERNEL_API_LOGDEST=debug.log
+
+python my_script.py
+```
+
+Output in `debug.log`:
+```
+================================================================================
+[2026-03-19 00:47:30] SGLang Kernel API Call: sglang.quant_method.UnquantizedLinearMethod.apply
+Positional input arguments:
+  arg[0]=QKVParallelLinear(
+      repr=QKVParallelLinear(in_features=1024, output_features=4096, bias=False, tp_size=1, gather_output=False)
+    )
+  arg[1]=Tensor(
+      shape=(1, 1024)
+      dtype=torch.bfloat16
+      device=cuda:0
+      requires_grad=False
+      is_contiguous=True
+    )
+  arg[2]=None
+Output:
+  return=Tensor(
+      shape=(1, 4096)
+      dtype=torch.bfloat16
+      device=cuda:0
+      requires_grad=False
+      is_contiguous=True
+    )
+```
+
+This is a real level-3 excerpt captured from `Qwen/Qwen3-0.6B`.
+
+### Full Logging (With Tensor Statistics)
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=5
+export SGLANG_KERNEL_API_LOGDEST=debug.log
+
+python my_script.py
+```
+
+Additional output:
+```
+================================================================================
+[2026-03-19 01:00:42] SGLang Kernel API Call: diffusion.quant_method.UnquantizedLinearMethod.apply
+Positional input arguments:
+  arg[1]=Tensor(
+      shape=(1, 77, 768)
+      dtype=torch.bfloat16
+      device=cuda:0
+      requires_grad=False
+      is_contiguous=True
+      min=-27.250000
+      max=28.500000
+      mean=0.011723
+      nan_count=0
+      inf_count=0
+    )
+Output:
+  return=Tensor(
+      shape=(1, 77, 2304)
+      dtype=torch.bfloat16
+      device=cuda:0
+      requires_grad=False
+      is_contiguous=True
+      min=-8.937500
+      max=9.375000
+      mean=0.009460
+      nan_count=0
+      inf_count=0
+    )
+```
+
+This is a real level-5 excerpt captured from `black-forest-labs/FLUX.1-dev`.
+
+### Crash-Safe Dumps (Inputs Saved Before Execution)
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=10
+export SGLANG_KERNEL_API_LOGDEST=debug.log
+export SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps
+
+python my_script.py
+```
+
+At level 10, SGLang saves the inputs before execution. If the kernel crashes, the dump directory still contains the inputs and exception metadata.
+
+If CUDA graph capture is active, tensor dumps are skipped automatically to avoid capture-time CUDA errors. In that case, you still get the kernel API call log, but not `inputs.pt` / `outputs.pt`.
+
+Level-10 dumps are best understood as crash-safe call snapshots. They always preserve the observed call boundary. They do not guarantee one-click replay for every method, because some methods depend on module state that is not serialized into the dump.
+
+Real level-10 dump layout from `Qwen/Qwen3-0.6B`:
+
+```text
+/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps
+/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001
+/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/inputs.pt
+/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/metadata.json
+/tmp/sglang_kernel_api_validation/qwen_qwen3_0_6b_level10_dumps/20260319_004821_182_pid919286_RotaryEmbedding.forward_call0001/outputs.pt
+```
+
+Real `metadata.json` excerpt:
+
+```json
+{
+  "function_name": "RotaryEmbedding.forward",
+  "timestamp": "20260319_004821_182",
+  "process_id": 919286,
+  "execution_status": "completed",
+  "input_tensor_keys": ["arg_0", "arg_1", "arg_2"],
+  "output_tensor_keys": ["result_0", "result_1"]
+}
+```
+
+## Step 2: Reproduce an LLM CUDA Crash
+
+Create a temporary reproducer:
+
+```bash
+python3 - <<'PY'
+from pathlib import Path
+Path("/tmp/sglang_llm_crash.py").write_text(
+    "import torch\\n"
+    "import torch.nn.functional as F\\n"
+    "from sglang.srt.utils.custom_op import register_custom_op\\n\\n"
+    "def _fake_embedding(indices, table):\\n"
+    "    return torch.empty((*indices.shape, table.shape[-1]), device=table.device, dtype=table.dtype)\\n\\n"
+    "@register_custom_op(op_name='mock_llm_cuda_crash', fake_impl=_fake_embedding)\\n"
+    "def mock_llm_cuda_crash(indices, table):\\n"
+    "    out = F.embedding(indices, table)\\n"
+    "    torch.cuda.synchronize()\\n"
+    "    return out\\n\\n"
+    "table = torch.randn(4, 8, device='cuda', dtype=torch.float16)\\n"
+    "indices = torch.tensor([0, 7], device='cuda', dtype=torch.long)\\n"
+    "mock_llm_cuda_crash(indices, table)\\n"
+)
+PY
+
+SGLANG_KERNEL_API_LOGLEVEL=1 \
+SGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level1.log \
+python3 /tmp/sglang_llm_crash.py
+```
+
+What to expect:
+- The script exits with a CUDA `device-side assert`
+- The log still contains the last API boundary before the crash
+
+Try the same example at level 3:
+
+```bash
+SGLANG_KERNEL_API_LOGLEVEL=3 \
+SGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level3.log \
+python3 /tmp/sglang_llm_crash.py
+```
+
+Now the log shows tensor metadata before the crash.
+
+Try level 10:
+
+```bash
+SGLANG_KERNEL_API_LOGLEVEL=10 \
+SGLANG_KERNEL_API_LOGDEST=/tmp/sglang_llm_level10.log \
+SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_llm_level10_dumps \
+python3 /tmp/sglang_llm_crash.py
+```
+
+Now you should see:
+- A log entry for `sglang.custom_op.mock_llm_cuda_crash`
+- A dump directory with `inputs.pt`
+- `metadata.json` showing `execution_status: "exception"`
+- No `outputs.pt`, because the kernel crashed before producing output
+
+For real-model success-path level-10 dumps, it is often easier to temporarily disable CUDA graph and piecewise CUDA graph for the debug run.
+
+## Step 3: Reproduce a Diffusion CUDA Crash
+
+Create a temporary diffusion-side reproducer:
+
+```bash
+python3 - <<'PY'
+from pathlib import Path
+Path("/tmp/sglang_diffusion_crash.py").write_text(
+    "import torch\\n"
+    "import torch.nn.functional as F\\n"
+    "from sglang.multimodal_gen.runtime.layers.utils import register_custom_op\\n\\n"
+    "def _fake_embedding(positions, cache):\\n"
+    "    return torch.empty((*positions.shape, cache.shape[-1]), device=cache.device, dtype=cache.dtype)\\n\\n"
+    "@register_custom_op(op_name='mock_diffusion_cuda_crash', fake_impl=_fake_embedding)\\n"
+    "def mock_diffusion_cuda_crash(positions, cache):\\n"
+    "    out = F.embedding(positions, cache)\\n"
+    "    torch.cuda.synchronize()\\n"
+    "    return out\\n\\n"
+    "cache = torch.randn(4, 64, device='cuda', dtype=torch.float16)\\n"
+    "positions = torch.tensor([0, 9], device='cuda', dtype=torch.long)\\n"
+    "mock_diffusion_cuda_crash(positions, cache)\\n"
+)
+PY
+
+SGLANG_KERNEL_API_LOGLEVEL=1 \
+SGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level1.log \
+python3 /tmp/sglang_diffusion_crash.py
+```
+
+Try level 3:
+
+```bash
+SGLANG_KERNEL_API_LOGLEVEL=3 \
+SGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level3.log \
+python3 /tmp/sglang_diffusion_crash.py
+```
+
+Try level 10:
+
+```bash
+SGLANG_KERNEL_API_LOGLEVEL=10 \
+SGLANG_KERNEL_API_LOGDEST=/tmp/sglang_diffusion_level10.log \
+SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_diffusion_level10_dumps \
+python3 /tmp/sglang_diffusion_crash.py
+```
+
+If your local environment has unrelated FlashInfer import issues, resolve them in the shell before running the example. The example itself does not set any `FLASHINFER_*` environment variable.
+
+## Step 4: Multi-Process Debugging
+
+When running with multiple GPUs or worker processes, use `%i` in the log path:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+export SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log
+
+torchrun --nproc_per_node=4 my_script.py
+```
+
+This creates separate logs such as:
+- `debug_rank_12345.log`
+- `debug_rank_12346.log`
+- `debug_rank_12347.log`
+- `debug_rank_12348.log`
+
+Real multi-process example from a 2-GPU `Qwen/Qwen2.5-0.5B-Instruct` run:
+
+```text
+/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950201.log
+/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950349.log
+/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950350.log
+/tmp/sglang_kernel_api_validation_multi/qwen_qwen2_5_0_5b_instruct_level3_950351.log
+```
+
+You should usually do the same for level-10 dump directories:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=10
+export SGLANG_KERNEL_API_LOGDEST=debug_rank_%i.log
+export SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps_%i
+```
+
+This avoids multiple ranks writing into the same dump directory tree.
+
+## Step 5: Filter Level-10 Dumps
+
+If level 10 is too noisy, restrict dumps to specific APIs:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=10
+export SGLANG_KERNEL_API_LOGDEST=debug.log
+export SGLANG_KERNEL_API_DUMP_DIR=/tmp/sglang_kernel_api_dumps
+export SGLANG_KERNEL_API_DUMP_INCLUDE='sglang.custom_op.*'
+export SGLANG_KERNEL_API_DUMP_EXCLUDE='*.fake_impl'
+```
+
+`SGLANG_KERNEL_API_DUMP_INCLUDE` and `SGLANG_KERNEL_API_DUMP_EXCLUDE` use shell-style wildcard matching.
+
+## Step 6: Common CUDA Errors and What to Check
+
+### Illegal Memory Access or Device-Side Assert
+
+**Typical errors**:
+```
+RuntimeError: CUDA error: an illegal memory access was encountered
+torch.AcceleratorError: CUDA error: device-side assert triggered
+```
+
+Use:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+```
+
+Check in the logs:
+- ✅ Tensor shapes
+- ✅ Tensor dtypes
+- ✅ CUDA vs CPU device placement
+- ✅ Tensor stride / contiguity
+- ✅ Whether the failing call has inputs logged but no outputs logged
+
+Typical shape-mismatch pattern:
+
+```text
+SGLang Kernel API Call: ...
+arg[0]=Tensor(shape=(..., 128), ...)   # ✅ expected dimension
+arg[1]=Tensor(shape=(..., 64), ...)    # ❌ mismatch
+```
+
+This often points to head-dim, hidden-dim, or cache-layout mismatch rather than a random CUDA failure.
+
+### NaN or Inf
+
+Use:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=5
+```
+
+Check:
+- `min`
+- `max`
+- `mean`
+- `nan_count`
+- `inf_count`
+
+Typical bad pattern:
+
+```text
+Tensor(
+  ...
+  min=-1234567.000000   # ❌ suspiciously large
+  max=9876543.000000    # ❌ suspiciously large
+  mean=nan              # ❌ bad
+  nan_count=128         # ❌ found NaNs
+  inf_count=0           # ✅ no Infs here
+)
+```
+
+This usually means the bad values were already present before the crashing kernel.
+
+### Out of Memory
+
+Use:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+```
+
+Check:
+- Unexpectedly large tensor shapes
+- Batch size
+- Sequence length
+- Frame count or image resolution in diffusion workloads
+
+Also check whether a supposedly per-token or per-frame tensor accidentally became full-sequence or full-image sized.
+
+Typical bad pattern:
+
+```text
+Tensor(
+  shape=(1024, 8192, 128, 128)   # ❌ way too large
+  ...
+)
+```
+
+### Example: Spot a Shape Bug from the Log
+
+Suppose the failing API log looks like this:
+
+```text
+[2026-03-19 00:47:30] SGLang Kernel API Call: RotaryEmbedding.forward
+Positional input arguments:
+  arg[0]=Tensor(shape=(1, 8), dtype=torch.int64, ...)
+  arg[1]=Tensor(shape=(1, 8, 8, 256), dtype=torch.bfloat16, ...)    # ✅ query
+  arg[2]=Tensor(shape=(1, 8, 4, 64), dtype=torch.bfloat16, ...)     # ❌ key head_dim mismatch
+```
+
+What this tells you:
+- ✅ positions look reasonable
+- ✅ query looks plausible
+- ❌ key last dimension is inconsistent with the expected rotary/head dimension
+
+That usually means the bug is in projection layout, head packing, or cache format rather than in the rotary kernel itself.
+
+## Step 7: Combine with compute-sanitizer
+
+For harder bugs, combine kernel API logging with CUDA memory checking:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+export SGLANG_KERNEL_API_LOGDEST=debug.log
+
+compute-sanitizer --tool memcheck python3 /tmp/sglang_llm_crash.py
+```
+
+Use `debug.log` to see the exact inputs that reached the crashing API boundary.
+
+Typical `compute-sanitizer` output:
+
+```text
+========= COMPUTE-SANITIZER
+========= Invalid __global__ write of size 4 bytes
+=========     at 0x1234 in SomeKernel
+=========     by thread (256,0,0) in block (10,0,0)
+=========     Address 0x... is out of bounds
+```
+
+Use the sanitizer output to identify the failing kernel and use `debug.log` to identify the exact tensors that reached the API boundary right before it.
+
+If you need more synchronous host-side error reporting, you can try `CUDA_LAUNCH_BLOCKING=1` as a separate follow-up experiment. It is not part of the default workflow because it changes execution timing and can hide concurrency-related behavior.
+
+## Step 8: Combine with cuda-gdb
+
+For crashes that need a stack trace instead of only memory diagnostics:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+export SGLANG_KERNEL_API_LOGDEST=debug.log
+
+cuda-gdb --args python3 /tmp/sglang_llm_crash.py
+```
+
+Inside `cuda-gdb`:
+
+```text
+(cuda-gdb) run
+(cuda-gdb) where
+```
+
+Then correlate the backtrace with `debug.log`.
+
+## Step 9: Kernel-Level Debugging with printf()
+
+When you own the CUDA kernel, `printf()` is still useful for narrowing down bad indices, bad launch geometry, or broken state propagation.
+
+Basic pattern:
+
+```cpp
+__global__ void MyKernel(const float* input, float* output, int n) {
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+  if (threadIdx.x == 0 && blockIdx.x == 0) {
+    printf("n=%d input0=%f\n", n, input[0]);
+  }
+
+  if (idx < n) {
+    output[idx] = input[idx] * 2.0f;
+  }
+}
+```
+
+After launch, force the output to flush:
+
+```python
+my_kernel(...)
+torch.cuda.synchronize()
+```
+
+For warp-specialized kernels, do not blindly print only on `threadIdx.x == 0`. Pick one representative thread per warp or per specialization group instead.
+
+### Warp-Specialized Kernels: Choosing the Right Print Thread
+
+Problem:
+- `threadIdx.x == 0` only prints from the first warp in the block
+- for warp-specialized kernels, that often misses the warp or group that is actually wrong
+
+Better pattern:
+
+```cpp
+__global__ void WarpSpecializedKernel(...) {
+  // Example: first lane of each warp
+  if ((threadIdx.x % 32) == 0) {
+    printf("warp=%d\n", threadIdx.x / 32);
+  }
+}
+```
+
+Or, if the kernel is organized in larger specialization groups, print once per group instead of once per block.
+
+Common mistake:
+
+```cpp
+// Only warp 0 prints
+if (threadIdx.x == 0) {
+  printf("warp=%d\n", threadIdx.x / 32);
+}
+```
+
+### Quick Reference
+
+| Kernel Type | Print Condition | Notes |
+|----------|----------|-------------|
+| Simple kernel | `threadIdx.x == 0` | One thread per block is usually enough |
+| Warp-specialized kernel | one representative lane per warp | e.g. `threadIdx.x % 32 == 0` |
+| Group-specialized kernel | one representative lane per group | choose based on the kernel's scheduling layout |
+
+### Other Kernel Debugging Tools
+
+```cpp
+assert(value >= 0.0f && "value must be non-negative");
+static_assert(BLOCK_SIZE % 32 == 0, "BLOCK_SIZE must be warp aligned");
+```
+
+## Environment Variables Reference
+
+| Variable | Values | Description |
+|----------|--------|-------------|
+| `SGLANG_KERNEL_API_LOGLEVEL` | `0` | No logging (default) |
+|  | `1` | Function names only |
+|  | `3` | Inputs and outputs with metadata |
+|  | `5` | Level 3 plus tensor statistics |
+|  | `10` | Level 5 plus crash-safe tensor dumps |
+| `SGLANG_KERNEL_API_LOGDEST` | `stdout` | Log to stdout |
+|  | `stderr` | Log to stderr |
+|  | `<path>` | Log to file |
+|  | `log_%i.txt` | `%i` expands to process ID |
+| `SGLANG_KERNEL_API_DUMP_DIR` | `<path>` | Directory for level-10 dumps |
+| `SGLANG_KERNEL_API_DUMP_INCLUDE` | wildcard list | Only dump matching API names |
+| `SGLANG_KERNEL_API_DUMP_EXCLUDE` | wildcard list | Skip matching API names |
+
+## Best Practices
+
+### 1. Start with Level 3
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+```
+
+Level 3 is usually enough to catch wrong shapes, wrong dtypes, and wrong devices.
+
+### 2. Use Level 5 for Numerical Issues
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=5
+```
+
+Use it when you suspect NaN or Inf values.
+
+### 3. Use Level 10 for Crash Reproduction
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=10
+```
+
+This is the most useful mode when the process crashes before you can inspect live tensors.
+
+If you need successful input/output dumps from a real model run, temporarily disable CUDA graph for that debug session.
+
+When level 10 is too noisy, pair it with `SGLANG_KERNEL_API_DUMP_INCLUDE` / `SGLANG_KERNEL_API_DUMP_EXCLUDE` instead of dumping every covered API.
+
+### 4. Log to File for Crashes
+
+```bash
+export SGLANG_KERNEL_API_LOGDEST=crash.log
+```
+
+File logs are safer than stdout when the process aborts.
+
+### 5. Disable Logging in Production
+
+```bash
+unset SGLANG_KERNEL_API_LOGLEVEL
+```
+
+When disabled, the decorator returns the original callable and adds no runtime logging overhead.
+
+## Troubleshooting
+
+### No Logs Appear
+
+Check:
+1. `echo $SGLANG_KERNEL_API_LOGLEVEL`
+2. `echo $SGLANG_KERNEL_API_LOGDEST`
+3. Whether the failing path goes through a covered API boundary
+
+### Too Much Output
+
+Reduce the level:
+
+```bash
+export SGLANG_KERNEL_API_LOGLEVEL=3
+```
+
+### Statistics Are Skipped During CUDA Graph Capture
+
+If you see:
+```text
+statistics=[skipped: CUDA graph capture in progress]
+```
+
+That is expected. Level-5 statistics are intentionally skipped during CUDA graph capture to avoid synchronization side effects.
+
+### Tensor Dumps Are Skipped During CUDA Graph Capture
+
+If you see:
+```text
+Tensor dump skipped: CUDA graph capture in progress
+```
+
+That is also expected. Level-10 dumps require copying tensors to CPU, which is not allowed during CUDA graph capture.

third_party/sglang/.claude/skills/generate-profile/SKILL.md

@@ -0,0 +1,141 @@
+---
+name: generate-profile
+description: Generate an e2e profiling trace of an SGLang server run. Launches a server, validates accuracy, captures a Chrome-compatible trace, and returns the profile path.
+---
+
+# Generate an E2E Profile of an SGLang Server Run
+
+This skill launches an SGLang server, validates it with a quick accuracy test, generates a profiling trace, and returns the profile file path.
+
+## Prerequisites
+
+- A working SGLang installation (`pip install -e .` or equivalent)
+- At least one available CUDA GPU
+
+## Step-by-step Workflow
+
+### Step 1: Launch the server
+
+```bash
+CUDA_VISIBLE_DEVICES=<gpu_id> sglang serve --model-path <model> --port <port> &
+```
+
+- Default model: `Qwen/Qwen3-8B` (good balance of speed and quality)
+- Default port: `30000`
+- The server runs in the background. Save the PID for cleanup.
+- Use the GPU specified by the user's preferences (check memory files for GPU preferences).
+
+### Step 2: Wait for server readiness
+
+Poll the health endpoint until the server is ready:
+
+```bash
+for i in $(seq 1 120); do
+  if curl -s http://127.0.0.1:<port>/health 2>/dev/null | grep -q "ok\|healthy"; then
+    echo "Server ready"
+    break
+  fi
+  sleep 5
+done
+```
+
+The server prints **"The server is fired up and ready to roll!"** to stdout when ready. The health endpoint returns 200 once the server can accept requests.
+
+Typical startup time: 30-90 seconds depending on model size and whether CUDA graphs are being compiled.
+
+### Step 3: Validate accuracy (sanity check)
+
+```bash
+python3 -m sglang.test.few_shot_gsm8k --num-q 20
+```
+
+- Expected accuracy: **> 0.8** for capable models (Qwen3-8B, Llama-3.1-8B-Instruct, etc.)
+- This is a quick sanity check, not a rigorous benchmark.
+- If accuracy is unexpectedly low, something is wrong — do not proceed to profiling.
+
+### Step 4: Generate the profile
+
+```bash
+python3 -m sglang.test.send_one --profile
+```
+
+This command:
+1. Sends a request to the server
+2. Triggers the profiler for 5 steps (default)
+3. Generates a trace file under `/tmp/<timestamp>/`
+4. The trace directory contains:
+   - `<timestamp>-TP-0.trace.json.gz` — Chrome trace format (open in `chrome://tracing` or Perfetto)
+   - `server_args.json` — the server configuration used
+
+**Output format:**
+```
+Dump profiling traces to /tmp/<timestamp>
+```
+
+The profile path is printed to stdout. Parse it from the output.
+
+**Optional flags:**
+- `--profile-steps N` — number of profiling steps (default: 5)
+- `--profile-by-stage` — profile by stage (prefill/decode separately)
+- `--profile-prefix <path>` — custom output prefix
+
+### Step 5: Kill the server
+
+```bash
+pkill -9 -f "sglang.launch_server\|sglang serve\|sglang.srt"
+```
+
+Wait a moment and verify no sglang processes remain:
+```bash
+sleep 2 && pgrep -af "sglang serve" || echo "Server killed"
+```
+
+### Step 6: Report the profile path
+
+Return the profile directory path (e.g., `/tmp/1773999986.4769795`) and list its contents so the user knows what files were generated.
+
+## Example Full Run
+
+```bash
+# 1. Launch server
+source cleanup/bin/activate
+CUDA_VISIBLE_DEVICES=1 sglang serve --model-path Qwen/Qwen3-8B --port 30000 &
+
+# 2. Wait for ready
+for i in $(seq 1 120); do
+  curl -s http://127.0.0.1:30000/health | grep -q "ok" && break
+  sleep 5
+done
+
+# 3. Accuracy check
+python3 -m sglang.test.few_shot_gsm8k --num-q 20
+# Expected: Accuracy > 0.8
+
+# 4. Profile
+python3 -m sglang.test.send_one --profile
+# Output: "Dump profiling traces to /tmp/1773999986.4769795"
+
+# 5. Cleanup
+pkill -9 -f "sglang.launch_server\|sglang serve\|sglang.srt"
+sleep 2
+
+# 6. Check output
+ls -la /tmp/1773999986.4769795/
+# 1773999986.4851577-TP-0.trace.json.gz  (Chrome trace)
+# server_args.json                        (server config)
+```
+
+## Customization
+
+- **Different port**: Pass `--port <port>` and use `--host 127.0.0.1 --port <port>` for test commands
+- **Multi-GPU**: Use `--tp <N>` for tensor parallelism; trace files will be generated per TP rank
+- **Longer profile**: Use `--profile-steps 10` for more steps in the trace
+- **Stage profiling**: Use `--profile-by-stage` to separate prefill and decode phases
+
+## Viewing the Profile
+
+Open the `.trace.json.gz` file in:
+- **Perfetto UI**: https://ui.perfetto.dev/ (drag and drop the file)
+- **Chrome tracing**: `chrome://tracing` (load the file)
+
+Both support the gzipped Chrome trace format natively.

third_party/sglang/.claude/skills/sglang-bisect-ci-regression/SKILL.md

@@ -0,0 +1,219 @@
+# SGLang Bisect CI Regression
+
+Investigate a consistently failing CI test to find the root cause - whether it's a code regression from a specific PR, a hardware/runner-specific issue, or an environment change. Optionally reproduce the failure on a remote GPU server.
+
+## Slash Command
+
+`/sglang-bisect-ci-regression <test_name_or_ci_url> [ssh_target] [docker_container]`
+
+## When to Use This Skill
+
+- A CI test is failing consistently on main (scheduled runs)
+- You need to find which PR introduced a regression
+- You suspect a runner-specific or GPU-specific issue
+- You want to reproduce a CI failure on a remote server
+
+## Arguments
+
+- **First argument (required)**: Test file name (e.g. `test_lora_tp.py`) or a GitHub Actions job URL
+- **Second argument (optional)**: SSH target for remote reproduction (e.g. `user@host`)
+- **Third argument (optional)**: Docker container name on the SSH target (e.g. `sglang_dev`)
+
+If SSH target and docker container are not provided, the skill will only perform the CI log analysis and bisection, without remote reproduction. **Ask the user** for these if reproduction is needed and they weren't provided.
+
+## Background: Scheduled CI Runs
+
+SGLang uses the `pr-test.yml` workflow with **scheduled runs** (cron-triggered) to periodically test the `main` branch. These runs are the primary data source for detecting regressions:
+
+- **Workflow**: `pr-test.yml` with `event: schedule`
+- **Branch**: `main`
+- **Dashboard**: https://github.com/sgl-project/sglang/actions/workflows/pr-test.yml?query=event%3Aschedule
+- **Frequency**: Runs multiple times daily, each pinned to the HEAD of `main` at trigger time
+- **Purpose**: Catches regressions that slip through PR-level CI (e.g., interaction bugs between merged PRs, hardware-specific issues)
+
+Always use these scheduled runs (not PR-triggered runs) when bisecting regressions on `main`. The `--event schedule` filter in `gh run list` ensures you only see these periodic main-branch runs.
+
+## Workflow
+
+### Phase 1: Extract the Failure Signature
+
+1. **Get the failing test details from CI logs.** If given a URL, fetch logs directly. If given a test name, find recent scheduled runs of `pr-test.yml` on `main` that failed:
+
+```bash
+# List recent scheduled runs targeting main (the primary source of truth for regressions)
+# These are cron-triggered runs visible at:
+# https://github.com/sgl-project/sglang/actions/workflows/pr-test.yml?query=event%3Aschedule
+gh run list --repo sgl-project/sglang --workflow="pr-test.yml" --event schedule --branch main --limit 20 --json databaseId,conclusion,createdAt,headSha
+
+# Find the job containing the test
+gh run view {RUN_ID} --repo sgl-project/sglang --json jobs --jq '.jobs[] | select(.conclusion == "failure") | {name, conclusion, databaseId}'
+
+# Get the failure details
+gh run view {RUN_ID} --repo sgl-project/sglang --job {JOB_ID} --log 2>&1 | grep -E -B 5 -A 30 "AssertionError|FAIL|Error|{TEST_NAME}"
+```
+
+2. **Record the failure signature:**
+   - Exact error message and assertion
+   - Affected test method name
+   - Model/config involved
+   - Numeric values (e.g., tolerance diffs, scores)
+   - Whether the failure is deterministic (same values across runs)
+
+### Phase 2: Temporal Bisection
+
+3. **Find the boundary between passing and failing runs.** Walk through the scheduled run history (from the `pr-test.yml` schedule runs on `main`) to identify:
+   - Last known PASSING run (sha + date)
+   - First known FAILING run (sha + date)
+
+```bash
+# For each scheduled run, check the specific partition/job status
+gh run view {RUN_ID} --repo sgl-project/sglang --json jobs --jq '.jobs[] | select(.name == "{JOB_NAME}") | {conclusion, databaseId}'
+
+# Verify a specific test passed or failed in a run
+gh run view {RUN_ID} --repo sgl-project/sglang --job {JOB_ID} --log 2>&1 | grep -E "{TEST_NAME}|PASSED|FAILED|logprobs mismatch" | head -10
+```
+
+4. **List commits between the boundary:**
+
+```bash
+git log --oneline {LAST_PASS_SHA}..{FIRST_FAIL_SHA}
+```
+
+5. **Filter for relevant commits** that touch files related to the failing test (model layers, kernels, test utilities, etc.):
+
+```bash
+git log --oneline {LAST_PASS_SHA}..{FIRST_FAIL_SHA} -- {relevant_paths}
+```
+
+### Phase 3: Runner/Hardware Analysis
+
+6. **Check if the failure is runner-specific.** Extract the runner identity from each failing and passing run:
+
+```bash
+# Get runner name and machine
+gh run view {RUN_ID} --repo sgl-project/sglang --job {JOB_ID} --log 2>&1 | grep -E "Runner name|Machine name" | head -5
+
+# Get GPU/driver info
+gh run view {RUN_ID} --repo sgl-project/sglang --job {JOB_ID} --log 2>&1 | grep -i -E "NVIDIA-SMI|Driver Version|CUDA Version" | head -5
+
+# Get package versions
+gh run view {RUN_ID} --repo sgl-project/sglang --job {JOB_ID} --log 2>&1 | grep -E "sgl.kernel.*==|flashinfer.*==" | head -5
+```
+
+7. **Correlate runners with pass/fail outcomes.** Build a table:
+
+| Run ID | Date | Runner | GPU Type | Driver | Result |
+|--------|------|--------|----------|--------|--------|
+
+If all failures map to a specific runner type/GPU and all passes map to another, the issue is **hardware-specific**, not a code regression.
+
+### Phase 4: Code Analysis
+
+8. **If a code regression is suspected** (failures not runner-specific), examine the candidate commits:
+   - Read the changed files
+   - Understand how the changes could affect the failing test
+   - Look for prefill-vs-decode differences, TP-specific paths, kernel changes
+
+9. **If a hardware issue is suspected**, analyze:
+   - Kernel compatibility (CUDA compute capability)
+   - Driver version differences
+   - All-reduce / NCCL behavior differences
+   - CUDA graph capture differences across GPU architectures
+
+### Phase 5: Remote Reproduction (Optional)
+
+Only if SSH target and docker container were provided.
+
+10. **Verify the remote environment:**
+
+```bash
+ssh {SSH_TARGET} "docker exec {CONTAINER} nvidia-smi --query-gpu=name,driver_version --format=csv"
+ssh {SSH_TARGET} "docker exec {CONTAINER} pip show sgl-kernel sglang flashinfer-python 2>&1 | grep -E 'Name:|Version:'"
+```
+
+11. **Ensure latest code is installed.** If the container is stale, update:
+
+```bash
+# Try fetching latest main
+ssh {SSH_TARGET} "docker exec {CONTAINER} bash -c 'cd /path/to/sglang && git fetch origin main && git checkout origin/main'"
+# Or download and install from tarball if git auth fails
+ssh {SSH_TARGET} "docker exec {CONTAINER} bash -c 'cd /tmp && curl -L https://github.com/sgl-project/sglang/archive/refs/heads/main.tar.gz | tar xz && cd sglang-main && pip install -e \"python[all]\"'"
+# Reinstall (after git fetch)
+ssh {SSH_TARGET} "docker exec {CONTAINER} bash -c 'cd /path/to/sglang && pip install -e \"python[all]\"'"
+# Install test dependencies if needed
+ssh {SSH_TARGET} "docker exec {CONTAINER} pip install peft rouge-score"
+```
+
+12. **Create a minimal reproduction script** that:
+    - Uses `if __name__ == '__main__'` with `mp.set_start_method("spawn")`
+    - Runs the specific failing test configuration
+    - Prints key metrics (diffs, scores, outputs)
+    - Exits with code 1 on failure
+
+13. **Copy and run the reproduction script:**
+
+```bash
+scp /tmp/repro_script.py {SSH_TARGET}:/tmp/
+ssh {SSH_TARGET} "docker cp /tmp/repro_script.py {CONTAINER}:/tmp/"
+ssh {SSH_TARGET} "docker exec -e CUDA_VISIBLE_DEVICES=0,1 {CONTAINER} python3 /tmp/repro_script.py"
+```
+
+14. **Run control experiments** to isolate the variable:
+    - If suspecting TP issue: run with TP=1 as control
+    - If suspecting GPU issue: compare same code on different GPU
+    - If suspecting a specific commit: test before/after that commit
+
+### Phase 6: Report
+
+15. **Produce a structured report:**
+
+```markdown
+## CI Regression Bisection Report
+
+### Failure Signature
+- **Test**: {test_file}::{test_method}
+- **Error**: {exact error message}
+- **Key metrics**: {numeric values}
+- **Deterministic**: Yes/No
+
+### Root Cause Classification
+One of:
+- **Code Regression**: PR #{number} introduced the bug
+- **Hardware-Specific**: Fails on {GPU_TYPE}, passes on others
+- **Environment Change**: New runner/driver/package version
+- **Pre-existing Flakiness**: Intermittent, not a new regression
+
+### Evidence
+| Condition | Result |
+|-----------|--------|
+| {condition1} | PASS/FAIL |
+| {condition2} | PASS/FAIL |
+
+### Timeline
+- {date}: Last known pass ({sha}, {runner})
+- {date}: First known fail ({sha}, {runner})
+- {date}: Confirmed reproduction on {server}
+
+### Recommended Fix
+- **Short-term**: {workaround}
+- **Long-term**: {proper fix}
+```
+
+## Key Patterns to Recognize
+
+| Pattern | Diagnosis |
+|---------|-----------|
+| Same SHA passes on runner A, fails on runner B | Hardware/runner-specific |
+| All runners fail after commit X | Code regression from commit X |
+| Intermittent - same runner sometimes passes/fails | Flaky test or race condition |
+| Prefill OK but decode fails | TP/all-reduce issue in decode path |
+| Works with TP=1, fails with TP>1 | Tensor parallelism bug |
+| Exact same numeric diff every time | Deterministic bug, not flakiness |
+
+## Important Notes
+
+- **Always check runner identity** before concluding it's a code regression. Many "consistent" failures are actually runner-specific.
+- **Test partition assignments change over time** as tests are added/removed. A test may move between partitions, landing on different runner types.
+- **H200 runners** use `/root/actions-runner/` path and machine names like `gpu-h200-worker-*`. Non-H200 runners use `/public_sglang_ci/runner-*` paths.
+- When running remote reproduction, use `run_in_background` for long-running tests and check output with `TaskOutput`.
+- Container environments may be stale - always verify package versions match CI before drawing conclusions.

third_party/sglang/.claude/skills/write-sglang-test/SKILL.md

@@ -0,0 +1,444 @@
+---
+name: write-sglang-test
+description: Guide for writing SGLang CI/UT tests. Covers CustomTestCase, CI registration, server fixtures, model selection, mock testing, and test placement. Always read test/README.md for the full CI layout, how to run tests, and extra tips. Use when creating new tests, adding CI test cases, writing unit tests, or when the user asks to add tests for SGLang features.
+---
+
+# Writing SGLang CI / UT Tests
+
+This skill covers **how to write and register tests**. For CI pipeline internals (stage ordering, fast-fail, gating, partitioning, debugging CI failures), see the [CI workflow guide](../ci-workflow-guide/SKILL.md).
+
+## Core Rules
+
+1. **Always use `CustomTestCase`** — never raw `unittest.TestCase`. It ensures `tearDownClass` runs even when `setUpClass` fails, preventing resource leaks in CI.
+2. **`tearDownClass` must be defensive** — use `hasattr`/null checks before accessing resources (e.g. `cls.process`) that `setUpClass` may not have finished allocating.
+3. **Place tests in `test/registered/<category>/`** — except JIT kernel tests and benchmarks, which live in `python/sglang/jit_kernel/tests/` and `python/sglang/jit_kernel/benchmark/` (nested subfolders are allowed)
+4. **Reuse server fixtures** — inherit from `DefaultServerBase` or write `setUpClass`/`tearDownClass` with `popen_launch_server`
+5. **Prefer mock over real server** — when testing logic that doesn't need a server / engine launch (middleware, request routing, config validation, argument parsing), use `unittest.mock.patch` / `MagicMock` and place tests in `test/registered/unit/`. Only launch a real server when the test genuinely needs inference results or server lifecycle behavior.
+
+JIT kernel exception:
+- If the task is adding or updating code under `python/sglang/jit_kernel/`, prefer the `add-jit-kernel` skill first.
+- JIT kernel correctness tests use `python/sglang/jit_kernel/tests/**/test_*.py`.
+- JIT kernel benchmarks use `python/sglang/jit_kernel/benchmark/**/bench_*.py`.
+- Those files are still executed by `test/run_suite.py`, but through dedicated kernel suites rather than `test/registered/`.
+
+---
+
+## Model & Backend Selection
+
+| Scenario | Model | CI Registration | Suite |
+|----------|-------|-----------------|-------|
+| **Unit tests** (no server / engine launch) | None | `register_cpu_ci` (prefer) or `register_cuda_ci` | `stage-a-test-cpu` or `stage-b-test-1-gpu-small` |
+| **Common / backend-independent** (middleware, abort, routing, config, arg parsing) | `DEFAULT_SMALL_MODEL_NAME_FOR_TEST` (1B) | `register_cuda_ci` only | `stage-b-test-1-gpu-small` |
+| **Model-agnostic functionality** (sampling, session, OpenAI API features) | `DEFAULT_SMALL_MODEL_NAME_FOR_TEST` (1B) | `register_cuda_ci` (+ AMD if relevant) | `stage-b-test-1-gpu-small` |
+| **General performance** (single node, no spec/DP/parallelism) | `DEFAULT_MODEL_NAME_FOR_TEST` (8B) | `register_cuda_ci` | `stage-b-test-1-gpu-large` |
+| **Bigger features** (spec, DP, TP, disaggregation) | Case by case | Case by case | See suite table below |
+
+**Key principle for E2E tests**: Do NOT add `register_amd_ci` unless the test specifically exercises AMD/ROCm code paths. Common E2E tests just need any GPU to run — duplicating across backends wastes CI time with no extra coverage.
+
+### All model constants
+
+Defined in `python/sglang/test/test_utils.py`:
+
+| Constant | Model | When to use |
+|----------|-------|-------------|
+| `DEFAULT_SMALL_MODEL_NAME_FOR_TEST` | Llama-3.2-1B-Instruct | Common features, model-agnostic tests |
+| `DEFAULT_SMALL_MODEL_NAME_FOR_TEST_BASE` | Llama-3.2-1B | Base (non-instruct) model tests |
+| `DEFAULT_MODEL_NAME_FOR_TEST` | Llama-3.1-8B-Instruct | General performance (single node) |
+| `DEFAULT_MOE_MODEL_NAME_FOR_TEST` | Mixtral-8x7B-Instruct | MoE-specific tests |
+| `DEFAULT_SMALL_EMBEDDING_MODEL_NAME_FOR_TEST` | — | Embedding tests |
+| `DEFAULT_SMALL_VLM_MODEL_NAME_FOR_TEST` | — | Vision-language tests |
+
+### Naming Conventions
+
+- **Suite**: `stage-{a,b,c}-test-{gpu_count}-gpu-{hardware}` (e.g., `stage-b-test-1-gpu-small`)
+- **CI runner**: `{gpu_count}-gpu-{hardware}` (e.g., `1-gpu-5090`, `4-gpu-h100`, `8-gpu-h200`)
+
+### All CI Suites
+
+#### Per-commit (CUDA)
+
+| Suite | Runner (label) | Description |
+|-------|----------------|-------------|
+| `stage-a-test-1-gpu-small` | `1-gpu-5090` | Quick checks on a small NVIDIA GPU before heavier stages |
+| `stage-a-test-cpu` | `ubuntu-latest` | CPU-only unit tests |
+| `stage-b-test-1-gpu-small` | `1-gpu-5090` | Core engine tests that fit a 5090-class card |
+| `stage-b-test-1-gpu-large` | `1-gpu-h100` | Tests that need H100-class memory or kernels (e.g. FA3) |
+| `stage-b-test-2-gpu-large` | `2-gpu-h100` | Two-GPU correctness and parallelism (TP/PP) on H100 |
+| `stage-b-test-4-gpu-b200` | `4-gpu-b200` | Early Blackwell coverage (SM100+ paths) on four GPUs |
+| `stage-b-kernel-unit-1-gpu-large` | `1-gpu-h100` | JIT kernel correctness tests under `python/sglang/jit_kernel/tests/` |
+| `stage-b-kernel-unit-8-gpu-h200` | `8-gpu-h200` | Multi-GPU JIT kernel correctness tests under `python/sglang/jit_kernel/tests/` |
+| `stage-b-kernel-benchmark-1-gpu-large` | `1-gpu-h100` | JIT kernel benchmark files under `python/sglang/jit_kernel/benchmark/` |
+| `stage-c-test-4-gpu-h100` | `4-gpu-h100` | Large 4-GPU H100 integration and scaling tests |
+| `stage-c-test-8-gpu-h200` | `8-gpu-h200` | Large 8-GPU H200 runs for big models and parallelism |
+| `stage-c-test-8-gpu-h20` | `8-gpu-h20` | Large 8-GPU H20 runs for big models |
+| `stage-c-test-deepep-4-gpu-h100` | `4-gpu-h100` | DeepEP expert-parallel and networking on four H100s |
+| `stage-c-test-deepep-8-gpu-h200` | `8-gpu-h200` | DeepEP at 8-GPU H200 scale |
+| `stage-c-test-8-gpu-b200` | `8-gpu-b200` | 8-GPU B200 suite (registered but not yet wired to a workflow) |
+| `stage-c-test-4-gpu-b200` | `4-gpu-b200` | 4-GPU B200 suite for large models on Blackwell |
+| `stage-c-test-4-gpu-gb200` | `4-gpu-gb200` | 4-GPU GB200 suite for large models on Grace Blackwell |
+
+#### Per-commit (AMD)
+
+| Suite | Runner (label) | Description |
+|-------|----------------|-------------|
+| `stage-a-test-1-gpu-small-amd` | `linux-mi325-1gpu-sglang` | Quick checks on one MI325-class GPU |
+| `stage-b-test-1-gpu-small-amd` | `linux-mi325-1gpu-sglang` | Core 1-GPU AMD tests (14 partitions) |
+| `stage-b-test-1-gpu-small-amd-nondeterministic` | `linux-mi325-1gpu-sglang` | Non-deterministic 1-GPU AMD tests |
+| `stage-b-test-1-gpu-small-amd-mi35x` | `linux-mi35x-gpu-1` | 1-GPU tests on MI35x hardware |
+| `stage-b-test-1-gpu-large-amd` | `linux-mi325-1gpu-sglang` | Large 1-GPU AMD tests (2 partitions) |
+| `stage-b-test-2-gpu-large-amd` | `linux-mi325-2gpu-sglang` | 2-GPU ROCm correctness and parallel setups |
+| `stage-b-test-large-8-gpu-35x-disaggregation-amd` | `linux-mi35x-gpu-8.fabric` | PD disaggregation and RDMA on 8×MI35x fabric |
+| `stage-c-test-4-gpu-amd` | `linux-mi325-4gpu-sglang` | 4-GPU AMD integration (2 partitions) |
+| `stage-c-test-large-8-gpu-amd` | `linux-mi325-8gpu-sglang` | 8-GPU MI325 scaling and integration |
+| `stage-c-test-large-8-gpu-amd-mi35x` | `linux-mi35x-gpu-8` | 8-GPU MI35x scaling (2 partitions) |
+
+
+### Per-commit (Ascend NPU)
+
+| Suite | Runner (label) | Description |
+| --- | --- | --- |
+| `per-commit-1-npu-a2` | `linux-aarch64-a2-1` | 1-NPU LLM CI machine |
+| `per-commit-2-npu-a2` | `linux-aarch64-a2-2` | 2-NPU LLM CI machine |
+| `per-commit-4-npu-a3` | `linux-aarch64-a3-4` | 4-NPU LLM CI machine |
+| `per-commit-16-npu-a3` | `linux-aarch64-a3-16` | 16-NPU LLM CI machine  |
+| `multimodal-gen-test-1-npu-a3` | `linux-aarch64-a3-2` | 1-NPU multimodal CI machine |
+| `multimodal-gen-test-2-npu-a3` | `linux-aarch64-a3-16` | 2-NPU multimodal CI machine |
+| `multimodal-gen-test-8-npu-a3` | `linux-aarch64-a3-16` | 8-NPU multimodal CI machine |
+
+#### Nightly
+
+Nightly suites are listed in `NIGHTLY_SUITES` in [`test/run_suite.py`](../../../test/run_suite.py). They run via `nightly-test-nvidia.yml`, `nightly-test-amd.yml` amd `nightly-test-npu.yml`, not `pr-test.yml`. Examples:
+
+- `nightly-1-gpu` (CUDA)
+- `nightly-kernel-1-gpu` (CUDA, JIT kernel full grids)
+- `nightly-kernel-8-gpu-h200` (CUDA, multi-GPU JIT kernel nightly)
+- `nightly-8-gpu-h200` (CUDA)
+- `nightly-eval-vlm-2-gpu` (CUDA)
+- `nightly-amd` (AMD)
+- `nightly-amd-8-gpu-mi35x` (AMD)
+- `nightly-1-npu-a3` (NPU)
+- `nightly-2-npu-a3` (NPU)
+- `nightly-4-npu-a3` (NPU)
+- `nightly-8-npu-a3` (NPU)
+- `nightly-16-npu-a3` (NPU)
+
+> **Note**: Multimodal diffusion uses `python/sglang/multimodal_gen/test/run_suite.py`, not `test/run_suite.py`.
+
+### Choosing a Suite
+
+Use the lightest suite that meets your test's needs:
+
+- **No GPU required** → `stage-a-test-cpu`
+- **Most small GPU tests** → `stage-b-test-1-gpu-small` (default choice)
+- **Need H100 memory or Hopper features** → `stage-b-test-1-gpu-large`
+- **JIT kernel correctness** → `stage-b-kernel-unit-1-gpu-large`
+- **JIT kernel benchmarks** → `stage-b-kernel-benchmark-1-gpu-large`
+- **Multi-GPU** → only when the test actually needs multiple GPUs
+
+---
+
+## Test File Templates
+
+### Unit Tests (no server / engine launch)
+
+See `test/registered/unit/README.md` for quick-start and rules. Unit tests live in `test/registered/unit/`, mirroring `python/sglang/srt/`:
+
+```python
+"""Unit tests for srt/<module>"""
+
+import unittest
+from unittest.mock import MagicMock, patch
+
+from sglang.srt.<module> import TargetClass
+from sglang.test.ci.ci_register import register_cpu_ci
+from sglang.test.test_utils import CustomTestCase
+
+register_cpu_ci(est_time=5, suite="stage-a-test-cpu")
+# Prefer CPU. Only use register_cuda_ci when the test truly needs a GPU.
+
+class TestTargetClass(CustomTestCase):
+    def test_basic_behavior(self):
+        obj = TargetClass(...)
+        self.assertEqual(obj.method(), expected)
+
+    @patch("sglang.srt.<module>.some_dependency")
+    def test_with_mock(self, mock_dep):
+        mock_dep.return_value = MagicMock()
+        # test logic with dependency mocked
+        ...
+
+
+if __name__ == "__main__":
+    unittest.main()
+```
+
+Use `unittest.mock.patch` / `MagicMock` to mock dependencies and isolate the logic under test. If the module transitively imports GPU-only packages (e.g. `sgl_kernel`), they can be stubbed so the test runs on CPU CI. See `test/registered/unit/README.md` for details and examples.
+
+**Quality bar** — test real logic (validation boundaries, state transitions, error paths, branching, etc.). Skip tests that just verify Python itself works (e.g., "does calling an abstract method raise `NotImplementedError`?", "does a dataclass store the field I assigned?"). Consolidate repetitive patterns into parameterized tests. No production code changes in test PRs.
+
+### E2E test (small model, server needed)
+
+```python
+import unittest
+
+import requests
+
+from sglang.srt.utils import kill_process_tree
+from sglang.test.ci.ci_register import register_cuda_ci
+from sglang.test.test_utils import (
+    DEFAULT_SMALL_MODEL_NAME_FOR_TEST,
+    DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH,
+    DEFAULT_URL_FOR_TEST,
+    CustomTestCase,
+    popen_launch_server,
+)
+
+register_cuda_ci(est_time=60, suite="stage-b-test-1-gpu-small")
+
+
+class TestMyFeature(CustomTestCase):
+    @classmethod
+    def setUpClass(cls):
+        cls.model = DEFAULT_SMALL_MODEL_NAME_FOR_TEST
+        cls.base_url = DEFAULT_URL_FOR_TEST
+        cls.process = popen_launch_server(
+            cls.model,
+            cls.base_url,
+            timeout=DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH,
+            other_args=["--arg1", "value1"],  # feature-specific args
+        )
+
+    @classmethod
+    def tearDownClass(cls):
+        if hasattr(cls, "process") and cls.process:
+            kill_process_tree(cls.process.pid)
+
+    def test_basic_functionality(self):
+        response = requests.post(
+            self.base_url + "/generate",
+            json={"text": "Hello", "sampling_params": {"max_new_tokens": 32}},
+        )
+        self.assertEqual(response.status_code, 200)
+
+
+if __name__ == "__main__":
+    unittest.main(verbosity=3)
+```
+
+### E2E test (8B model, server needed, performance)
+
+```python
+import time
+import unittest
+
+import requests
+
+from sglang.srt.utils import kill_process_tree
+from sglang.test.ci.ci_register import register_cuda_ci
+from sglang.test.test_utils import (
+    DEFAULT_MODEL_NAME_FOR_TEST,
+    DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH,
+    DEFAULT_URL_FOR_TEST,
+    CustomTestCase,
+    popen_launch_server,
+)
+
+register_cuda_ci(est_time=300, suite="stage-b-test-1-gpu-large")
+
+
+class TestMyFeaturePerf(CustomTestCase):
+    @classmethod
+    def setUpClass(cls):
+        cls.model = DEFAULT_MODEL_NAME_FOR_TEST
+        cls.base_url = DEFAULT_URL_FOR_TEST
+        cls.process = popen_launch_server(
+            cls.model,
+            cls.base_url,
+            timeout=DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH,
+        )
+
+    @classmethod
+    def tearDownClass(cls):
+        if hasattr(cls, "process") and cls.process:
+            kill_process_tree(cls.process.pid)
+
+    def test_latency(self):
+        start = time.perf_counter()
+        response = requests.post(
+            self.base_url + "/generate",
+            json={"text": "Hello", "sampling_params": {"max_new_tokens": 128}},
+        )
+        elapsed = time.perf_counter() - start
+        self.assertEqual(response.status_code, 200)
+        self.assertLess(elapsed, 5.0, "Latency exceeded threshold")
+
+
+if __name__ == "__main__":
+    unittest.main(verbosity=3)
+```
+
+---
+
+## Server Fixture Reuse
+
+For tests that only need a standard server, inherit from `DefaultServerBase` and override class attributes:
+
+```python
+from sglang.test.server_fixtures.default_fixture import DefaultServerBase
+
+class TestMyFeature(DefaultServerBase):
+    model = DEFAULT_SMALL_MODEL_NAME_FOR_TEST
+    other_args = ["--enable-my-feature"]
+
+    def test_something(self):
+        ...
+```
+
+Available fixtures in `python/sglang/test/server_fixtures/`:
+
+| Fixture | Use case |
+|---------|----------|
+| `DefaultServerBase` | Standard single-server tests |
+| `EagleServerBase` | EAGLE speculative decoding |
+| `PDDisaggregationServerBase` | Disaggregated prefill/decode |
+| `MMMUServerBase` | Multimodal VLM tests |
+
+---
+
+## CI Registration
+
+Every CI-discovered test file must call a registration function at module level:
+
+```python
+from sglang.test.ci.ci_register import (
+    register_cuda_ci,
+    register_amd_ci,
+    register_cpu_ci,
+    register_npu_ci,
+)
+
+# Per-commit test (small 1-gpu, runs on 5090)
+register_cuda_ci(est_time=80, suite="stage-b-test-1-gpu-small")
+
+# Per-commit test (large 1-gpu, runs on H100)
+register_cuda_ci(est_time=120, suite="stage-b-test-1-gpu-large")
+
+# Nightly-only test
+register_cuda_ci(est_time=200, suite="nightly-1-gpu", nightly=True)
+
+# Multi-backend test (only when testing backend-specific code paths)
+register_cuda_ci(est_time=80, suite="stage-a-test-1-gpu-small")
+register_amd_ci(est_time=120, suite="stage-a-test-1-gpu-small-amd")
+register_npu_ci(est_time=400, suite="nightly-8-npu-a3", nightly=True)
+
+# Temporarily disabled test
+register_cuda_ci(est_time=80, suite="stage-b-test-1-gpu-small", disabled="flaky - see #12345")
+```
+
+Parameters:
+- `est_time`: estimated runtime in seconds (used for CI partitioning)
+- `suite`: which CI suite to run in (see suite tables above)
+- `nightly=True`: for nightly-only tests (default `False` = per-commit)
+- `disabled="reason"`: temporarily disable with explanation
+
+**Key principle**: Only add `register_amd_ci` / `register_npu_ci` when the test exercises backend-specific code paths. Common E2E tests just need `register_cuda_ci` — duplicating across backends wastes CI time.
+
+### JIT Kernel Registration
+
+JIT kernel files live outside `test/registered/` but still use registration:
+
+```python
+from sglang.test.ci.ci_register import register_cuda_ci
+
+# Correctness tests in python/sglang/jit_kernel/tests/
+register_cuda_ci(est_time=30, suite="stage-b-kernel-unit-1-gpu-large")
+register_cuda_ci(est_time=120, suite="stage-b-kernel-unit-8-gpu-h200")
+
+# Benchmarks in python/sglang/jit_kernel/benchmark/
+register_cuda_ci(est_time=6, suite="stage-b-kernel-benchmark-1-gpu-large")
+
+# Optional nightly registration
+register_cuda_ci(est_time=120, suite="nightly-kernel-1-gpu", nightly=True)
+register_cuda_ci(est_time=120, suite="nightly-kernel-8-gpu-h200", nightly=True)
+```
+
+Keep `est_time` and `suite` as **literal values** — `run_suite.py` collects them by AST parsing
+
+---
+
+## Test Placement
+
+```
+test/
+├── registered/          # CI tests (auto-discovered by run_suite.py)
+│   ├── unit/            # No server / engine launch (see test/registered/unit/README.md)
+│   ├── kernels/         # CUDA kernel correctness (no server, GPU required)
+│   ├── sampling/        # test_penalty.py, test_sampling_params.py ...
+│   ├── sessions/        # test_session_control.py ...
+│   ├── openai_server/   # basic/, features/, validation/ ...
+│   ├── spec/            # eagle/, utils/ ...
+│   ├── models/          # model-specific accuracy tests
+│   ├── perf/            # performance benchmarks
+│   └── <category>/      # create new category if needed
+├── manual/              # Non-CI: debugging, one-off, manual verification
+└── run_suite.py         # CI runner (scans registered/ plus jit_kernel test/benchmark files)
+
+python/sglang/jit_kernel/
+├── tests/               # JIT kernel correctness tests (CI-discovered by test/run_suite.py)
+└── benchmark/           # JIT kernel benchmarks (CI-discovered by test/run_suite.py)
+```
+
+**Decision rule** (see also `test/registered/README.md`):
+- Component logic, no server → `registered/unit/`
+- JIT kernel correctness / benchmarks → `python/sglang/jit_kernel/tests/` or `python/sglang/jit_kernel/benchmark/`
+- Other kernel correctness → `registered/kernels/`
+- Server needed → `registered/<category>/`
+- Local debugging → `manual/`
+
+---
+
+## Eval Accuracy Mixins
+
+**Design philosophy**: Most test files don't care about eval logic — they only need a "does this feature break model output quality?" sanity check. The mixin pattern separates **what to test** (threshold) from **how to test** (run_eval, assertions, CI summary). Test classes declare thresholds as class attributes; the mixin provides the `test_*` method. Override when you need extra assertions (e.g. EAGLE accept length).
+
+Available mixins in `python/sglang/test/kits/eval_accuracy_kit.py`: `MMLUMixin`, `HumanEvalMixin`, `MGSMEnMixin`, `GSM8KMixin`. Can be combined freely. Read the source for attrs and defaults.
+
+```python
+class TestMyFeature(CustomTestCase, MMLUMixin):
+    mmlu_score_threshold = 0.65
+    mmlu_num_examples = 64
+    mmlu_num_threads = 32
+    # test_mmlu is inherited — no code needed
+```
+
+---
+
+## Key Utilities
+
+```python
+from sglang.test.test_utils import (
+    CustomTestCase,              # base class with retry logic
+    popen_launch_server,         # launch server subprocess
+    DEFAULT_URL_FOR_TEST,        # auto-configured base URL
+    DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH,  # 600s default
+    run_bench_serving,           # benchmark helper (launch + bench)
+)
+from sglang.srt.utils import kill_process_tree  # cleanup server
+```
+
+---
+
+## Checklist
+
+Before submitting a test:
+
+- [ ] Inherits from `CustomTestCase` (not `unittest.TestCase`)
+- [ ] Has `register_*_ci(...)` call at module level
+- [ ] Placed in `test/registered/<category>/`, unless this is a JIT kernel test/benchmark
+- [ ] JIT kernel work: files live in `python/sglang/jit_kernel/tests/` or `python/sglang/jit_kernel/benchmark/`
+- [ ] Backend-independent tests: `register_cuda_ci` only + smallest model
+- [ ] Logic that doesn't need a server / engine launch → unit test in `registered/unit/` (see Unit Tests section)
+- [ ] `setUpClass` launches server, `tearDownClass` kills it (if server-based)
+- [ ] `tearDownClass` is defensive — uses `hasattr`/null checks before accessing resources that may not have been allocated
+- [ ] Has `if __name__ == "__main__": unittest.main()`
+- [ ] `est_time` is reasonable (measure locally)

third_party/sglang/.codespellrc

@@ -0,0 +1,3 @@
+[codespell]
+ignore-words-list = ans, als, hel, boostrap, childs, te, vas, hsa, ment, cann, thi, makro, wil, rouge, PRIS
+skip = *.json,*.jsonl,*.patch,*.txt

third_party/sglang/.coveragerc

@@ -0,0 +1,16 @@
+[run]
+source = python/sglang/srt
+omit =
+    */test/*
+    */__pycache__/*
+
+[report]
+show_missing = true
+exclude_lines =
+    pragma: no cover
+    if __name__ == .__main__.:
+    raise NotImplementedError
+    if TYPE_CHECKING
+
+[html]
+directory = htmlcov

third_party/sglang/.devcontainer/Dockerfile

@@ -0,0 +1,35 @@
+FROM lmsysorg/sglang:dev
+
+# Create non-root user with specified UID and GID
+# NOTE: Replace with your own UID and GID. This is a workaround from https://github.com/microsoft/vscode-remote-release/issues/49#issuecomment-489060908.
+ARG HOST_UID=1003
+ARG HOST_GID=1003
+RUN groupadd -g $HOST_GID devuser && \
+    useradd -m -u $HOST_UID -g $HOST_GID -s /bin/zsh devuser
+
+# Give devuser sudo access
+RUN apt-get update && apt-get install -y sudo && \
+    echo "devuser ALL=(ALL) NOPASSWD:ALL" > /etc/sudoers.d/devuser && \
+    rm -rf /var/lib/apt/lists/* && \
+    apt-get clean
+
+# Set up oh-my-zsh for devuser
+RUN cp -r /root/.oh-my-zsh /home/devuser/.oh-my-zsh && \
+    cp /root/.zshrc /home/devuser/.zshrc && \
+    cp /root/.vimrc /home/devuser/.vimrc && \
+    cp /root/.tmux.conf /home/devuser/.tmux.conf && \
+    sed -i 's|/root/.oh-my-zsh|/home/devuser/.oh-my-zsh|g' /home/devuser/.zshrc && \
+    chown -R devuser:devuser /home/devuser/
+
+# Set workspace directory and ownership
+WORKDIR /sgl-workspace/sglang
+RUN chown -R devuser:devuser /sgl-workspace
+
+# Switch to devuser
+USER devuser
+
+# Install uv
+RUN curl -LsSf https://astral.sh/uv/install.sh | sh
+
+# Install rust
+RUN curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y

third_party/sglang/.devcontainer/devcontainer.json

@@ -0,0 +1,30 @@
+{
+    "name": "sglang",
+    "build": {
+        "dockerfile": "Dockerfile"
+    },
+    "remoteUser": "devuser",
+    "customizations": {
+        "vscode": {
+            "extensions": [
+                // Python development
+                "ms-python.python",
+                "charliermarsh.ruff",
+                // Rust development
+                "rust-lang.rust-analyzer",
+                "tamasfe.even-better-toml"
+            ]
+        }
+    },
+    "forwardPorts": [],
+    "runArgs": [
+        "--gpus",
+        "all"
+    ],
+    // The two lines below ensures that your local changes in the sglang
+    // repo is automatically synced to the sglang pip package installed
+    // in the dev docker container. You can remove / comment out these
+    // two lines if you prefer to sync code changes manually.
+    "workspaceMount": "source=${localWorkspaceFolder},target=/sgl-workspace/sglang,type=bind",
+    "workspaceFolder": "/sgl-workspace/sglang"
+}

third_party/sglang/.dockerignore

@@ -0,0 +1 @@
+.gitignore

third_party/sglang/.github/CODEOWNERS

@@ -0,0 +1,74 @@
+.github @merrymercy @Fridge003 @ispobock @Kangyan-Zhou @bingxche
+/docker @Fridge003 @ispobock @HaiShaw @ishandhanani @yctseng0211
+/docker/npu.Dockerfile @ping1jing2 @iforgetmyname
+/python/pyproject.toml @merrymercy @Fridge003 @ispobock
+/python/sglang/jit_kernel @DarkSharpness @BBuf @celve @HydraQYH @yuan-luo
+/python/sglang/jit_kernel/diffusion @yingluosanqian @BBuf @mickqian
+/python/sglang/multimodal_gen @mickqian @yhyang201 @ping1jing2
+/python/sglang/multimodal_gen/runtime/cache @DefTruth
+/python/sglang/multimodal_gen/runtime/layers @mickqian @yhyang201 @BBuf @yingluosanqian @ping1jing2
+/python/sglang/multimodal_gen/runtime/models/dits @mickqian @yhyang201 @BBuf @yingluosanqian @ping1jing2
+/python/sglang/srt/batch_invariant_ops @Fridge003 @hebiao064
+/python/sglang/srt/compilation @hebiao064 @Oasis-Git
+/python/sglang/srt/constrained @hnyls2002 @DarkSharpness
+/python/sglang/srt/disaggregation @ByronHsu @hnyls2002 @ShangmingCai
+/python/sglang/srt/disaggregation/ascend @ping1jing2 @iforgetmyname
+/python/sglang/srt/distributed @yizhang2077 @merrymercy @ch-wan
+/python/sglang/srt/distributed/device_communicators/mooncake_transfer_engine.py @ShangmingCai @stmatengss
+/python/sglang/srt/dllm @ClawSeven @btw616
+/python/sglang/srt/entrypoints @ispobock @CatherineSue @slin1237 @merrymercy @JustinTong0323
+/python/sglang/srt/entrypoints/engine_score_mixin.py @sundar24295s @chanh @fortunecookiee
+/python/sglang/srt/entrypoints/grpc_server.py @CatherineSue @slin1237
+/python/sglang/srt/entrypoints/openai/serving_score.py @sundar24295s @chanh @fortunecookiee
+/python/sglang/srt/eplb @fzyzcjy @ch-wan
+/python/sglang/srt/function_call @CatherineSue @JustinTong0323
+/python/sglang/srt/grpc @CatherineSue @slin1237
+/python/sglang/srt/hardware_backend/npu @ping1jing2 @iforgetmyname
+/python/sglang/srt/hardware_backend/npu/quantization @OrangeRedeng @TamirBaydasov @iforgetmyname
+/python/sglang/srt/layers @merrymercy @Ying1123 @Fridge003 @ispobock @HaiShaw @ch-wan @BBuf @Edwardf0t1
+/python/sglang/srt/layers/attention @merrymercy @Fridge003 @ispobock @Qiaolin-Yu @hebiao064 @HaiShaw
+/python/sglang/srt/layers/attention/fla @yizhang2077 @hebiao064 @yuan-luo
+/python/sglang/srt/layers/attention/hybrid_linear_attn_backend.py @yizhang2077 @hebiao064 @hanming-lu @yuan-luo
+/python/sglang/srt/layers/attention/mamba @yizhang2077 @hebiao064
+/python/sglang/srt/layers/attention/nsa @1am9trash @hubertlu-tw @kkHuang-amd @HaiShaw @Fridge003 @hlu1 @rainj-me
+/python/sglang/srt/layers/attention/vision.py @mickqian @yuan-luo @yhyang201
+/python/sglang/srt/layers/quantization @ch-wan @BBuf @Edwardf0t1 @FlamingoPg @AniZpZ @HaiShaw @b8zhong
+/python/sglang/srt/layers/quantization/quark @kkHuang-amd @yichiche @hubertlu-tw @1am9trash @BowenBao
+/python/sglang/srt/lora @Ying1123 @Fridge003 @lifuhuang @yushengsu-thu
+/python/sglang/srt/managers @merrymercy @Ying1123 @hnyls2002 @xiezhq-hermann
+/python/sglang/srt/managers/scheduler_pp_mixin.py @ShangmingCai @XucSh
+/python/sglang/srt/managers/tokenizer_manager_score_mixin.py @sundar24295s @chanh @fortunecookiee
+/python/sglang/srt/mem_cache @merrymercy @Ying1123 @hnyls2002 @xiezhq-hermann @hanming-lu @yizhang2077 @hzh0425 @ispobock
+/python/sglang/srt/model_executor @merrymercy @Ying1123 @hnyls2002 @Fridge003 @ispobock
+/python/sglang/srt/model_executor/piecewise_cuda_graph_runner.py @hebiao064
+/python/sglang/srt/models/deepseek_common @Fridge003 @ispobock @fzyzcjy @ch-wan
+/python/sglang/srt/models/deepseek_v2.py @fzyzcjy @zhyncs @ispobock @ch-wan @merrymercy @Fridge003
+/python/sglang/srt/models/transformers.py @adarshxs
+/python/sglang/srt/multimodal @mickqian @JustinTong0323 @yhyang201 @yuan-luo
+/python/sglang/srt/observability @merrymercy @fzyzcjy @sufeng-buaa
+/python/sglang/srt/ray @Qiaolin-Yu @xyuzh
+/python/sglang/srt/speculative @Ying1123 @merrymercy @hnyls2002
+/sgl-kernel @ispobock @BBuf @yizhang2077 @merrymercy @FlamingoPg @HaiShaw
+/sgl-model-gateway @slin1237 @CatherineSue
+/sgl-model-gateway/benches @slin1237
+/sgl-model-gateway/bindings/python @CatherineSue @key4ng @slin1237
+/sgl-model-gateway/e2e_test @CatherineSue @key4ng
+/sgl-model-gateway/examples/wasm @slin1237
+/sgl-model-gateway/src/config @slin1237
+/sgl-model-gateway/src/core @slin1237
+/sgl-model-gateway/src/data_connector @key4ng
+/sgl-model-gateway/src/grpc_client @CatherineSue @slin1237
+/sgl-model-gateway/src/mcp @key4ng @slin1237
+/sgl-model-gateway/src/policies @slin1237 @ByronHsu
+/sgl-model-gateway/src/proto @CatherineSue @slin1237
+/sgl-model-gateway/src/protocols @CatherineSue @key4ng
+/sgl-model-gateway/src/reasoning_parser @CatherineSue
+/sgl-model-gateway/src/routers @CatherineSue @key4ng @slin1237
+/sgl-model-gateway/src/tokenizer @slin1237 @CatherineSue
+/sgl-model-gateway/src/tool_parser @slin1237 @CatherineSue
+/sgl-model-gateway/src/wasm @slin1237
+/sgl-model-gateway/examples/wasm @slin1237
+/test/registered/core/test_score_api.py @sundar24295s @chanh @fortunecookiee
+/benchmark/prefill_only/bench_score.py @sundar24295s @chanh @fortunecookiee
+/test/srt/ascend @ping1jing2 @iforgetmyname
+/test/srt/test_modelopt* @Edwardf0t1

third_party/sglang/.github/FOLDER_README.md

@@ -0,0 +1,12 @@
+# Maintenance Tools
+
+This folder contains tools and workflows for automating maintenance tasks.
+
+## CI Permissions
+
+`CI_PERMISSIONS.json` defines the CI permissions granted to each user.
+Maintainers can directly edit the file to add entries with `"reason": "custom override"`.
+Maintainers can also run `update_ci_permission.py` to update it with some auto rules (e.g., top contributors in the last 90 days get full permissions).
+
+## Others
+- `MAINTAINER.md` defines the code maintenance model.

third_party/sglang/.github/ISSUE_TEMPLATE/1-bug-report.yml

@@ -0,0 +1,35 @@
+name: 🐞 Bug report
+description: Report a bug to help us reproduce and fix it.
+title: "[Bug] "
+labels: ['Bug']
+
+body:
+- type: checkboxes
+  attributes:
+    label: Checklist
+    options:
+      - label: I searched related issues but found no solution.
+      - label: The bug persists in the latest version.
+      - label: Issues without environment info and a minimal reproducible demo are hard to resolve and may receive no feedback.
+      - label: If this is not a bug report but a general question, please start a discussion at https://github.com/sgl-project/sglang/discussions. Otherwise, it will be closed.
+      - label: Please use English. Otherwise, it will be closed.
+- type: textarea
+  attributes:
+    label: Describe the bug
+    description: A clear, concise description of the bug.
+  validations:
+    required: true
+- type: textarea
+  attributes:
+    label: Reproduction
+    description: Command/script run and model used.
+    placeholder: Paste the command here.
+  validations:
+    required: true
+- type: textarea
+  attributes:
+    label: Environment
+    description: Run `python3 -m sglang.check_env` and paste output here. Issues without this will be closed.
+    placeholder: Paste environment output here.
+  validations:
+    required: true

third_party/sglang/.github/ISSUE_TEMPLATE/2-feature-request.yml

@@ -0,0 +1,23 @@
+name: 🚀 Feature request
+description: Suggest an idea for this project
+title: "[Feature] "
+
+body:
+- type: checkboxes
+  attributes:
+    label: Checklist
+    options:
+      - label: If this is not a feature request but a general question, please start a discussion at https://github.com/sgl-project/sglang/discussions. Otherwise, it will be closed.
+      - label: Please use English. Otherwise, it will be closed.
+- type: textarea
+  attributes:
+    label: Motivation
+    description: |
+      Clearly and concisely describe the feature's motivation.
+  validations:
+    required: true
+- type: textarea
+  attributes:
+    label: Related resources
+    description: |
+      Provide official releases or third-party implementations if available.

third_party/sglang/.github/MAINTAINER.md

@@ -0,0 +1,154 @@
+# SGLang Code Maintenance Model
+This document describes the code maintenance model for the SGLang project.
+Since SGLang is a large project involving multiple organizations and hardware platforms, we designed this model with the following goals:
+- Ensure a responsive and smooth review process.
+- Allow for fast iteration, so maintainers can sometimes bypass flaky CI tests for important PRs.
+
+## Role Descriptions
+There are four roles in this maintenance model. Some are custom roles, while others are predefined by GitHub.
+
+- **Merge Oncall**: The person who drives the PR merge process. They have strong area-specific expertise and uphold a high bar for code quality.
+  - Permission: Merge PRs. Bypass branch protection rules if needed.
+  - Responsibility: Shepherd the merge of PRs assigned to their area. Revert or hotfix any issues related to their merge (especially if they bypass).
+- **Codeowner**: The person who protects critical code. Without a bypass, each PR needs at least one Codeowner approval for each modified file protected by [CODEOWNERS](./CODEOWNERS). Please note that this role is not an honor but a significant responsibility because PRs cannot be merged without your approval (except when bypassed by a Merge Oncall).
+  - Permission: Approve PRs, allowing them to be merged without a bypass.
+  - Responsibility: Review PRs in a timely manner.
+- **Write**: A person with write permission to the SGLang repo.
+  - Permission: Merge PRs if they have passed required tests and been approved by Codeowners. This role cannot bypass branch protection rules.
+  - Responsibility: Review and merge PRs in a timely manner.
+- **CI Oncall**: A person who manages CI runners for specific hardware platforms.
+  - Permission: Add CI runners.
+  - Responsibility: Keep the CI runners up and running.
+
+__Note__: Difference between Merge Oncall and Codeowner
+- The Merge Oncall is an active role held by someone who actively tries to help merge PRs and can bypass CI if needed.
+- The Codeowner is a passive protection role provided by GitHub; it prevents accidental changes to critical code.
+- The list of Merge Oncalls is attached below. The list of Codeowners is in the [CODEOWNERS](./CODEOWNERS) file.
+
+__Note__: The permissions to trigger CI tests are defined separately according to these [rules](https://docs.sglang.io/developer_guide/contribution_guide.html#how-to-trigger-ci-tests).
+
+
+## Pull Request Merge Process
+1. The author submits a pull request (PR) and fills out the PR checklist.
+2. A bot assigns this PR to a Merge Oncall and @-mentions them. At the same time, GitHub will automatically request reviews from Codeowners.
+3. Someone tags the PR with a `run-ci` label ([help](https://docs.sglang.io/developer_guide/contribution_guide.html#how-to-trigger-ci-tests)). Then the author can trigger CI by pushing new commits.
+4. The Merge Oncall coordinates the review (e.g., asking people to review) and approves the PR; the Codeowners also approve the PR. If the assigned Merge Oncall is not responsive, the author can ping other related Merge Oncalls and Reviewers in the list below.
+5. The code can now be merged:
+   - **Ideal case:** For each modified file, one Codeowner has approved the PR. The PR has also passed the required CI tests. Then, anyone with write permission can merge the PR.
+   - **Exception:** In cases where it is difficult to meet all requirements (due to flaky CI or slow responses), a Merge Oncall can bypass branch protection to merge the PR.
+
+If you meet any issues during the merge, you can discuss in [slack channels](https://slack.sglang.io/): #pull-request, #ci-cd-build-release, #dev.
+
+## The List of Merge Oncalls and Reviewers
+This section lists the oncalls for each module or feature.
+The format is @github-username (Slack username).
+
+### Scheduler
+[@merrymercy](https://github.com/merrymercy) (Lianmin Zheng), [@hnyls2002](https://github.com/hnyls2002) (Liangsheng Yin), [@cctry](https://github.com/cctry) (Shiyang Chen)
+
+related files
+- python/sglang/srt/managers
+- python/sglang/srt/model_executor
+
+### Diffusion
+[@mickqian](https://github.com/mickqian) (Mick), [@BBuf](https://github.com/BBuf) (BBuf)
+
+related files
+- python/sglang/multimodal_gen
+
+### PD disaggregation
+[@ByronHsu](https://github.com/ByronHsu) (Byron Hsu), [@cctry](https://github.com/cctry) (Shiyang Chen), [@ShangmingCai](https://github.com/ShangmingCai) (Shangming Cai)
+
+related files
+- python/sglang/srt/disaggregation
+
+### KV Cache
+[@ispobock](https://github.com/ispobock) (Ke Bao), [@xiezhq-hermann](https://github.com/xiezhq-hermann) (Zhiqiang Xie)
+
+related files
+- python/sglang/srt/mem_cache
+
+### Parallelism
+[@ch-wan](https://github.com/ch-wan) (Cheng Wan), [@fzyzcjy](https://github.com/fzyzcjy) (Tom)
+
+related files
+- python/sglang/srt/eplb
+- python/sglang/srt/distributed
+- python/sglang/srt/layers/dp_attention.py
+
+### Kernel
+[@BBuf](https://github.com/BBuf) (BBuf)
+
+related files
+- python/sglang/jit_kernel
+- sgl-kernel
+
+### Speculative decoding
+[@hnyls2002](https://github.com/hnyls2002) (Liangsheng Yin), [@Qiaolin-Yu](https://github.com/Qiaolin-Yu) (Qiaolin Yu)
+
+related files
+- python/sglang/srt/speculative
+
+### NV and model-specific optimizations
+[@Fridge003](https://github.com/Fridge003) (Baizhou Zhang), [@ishandhanani](https://github.com/ishandhanani) (Ishan Dhanani), [@Qiaolin-Yu](https://github.com/Qiaolin-Yu) (Qiaolin Yu)
+
+related files
+- python/sglang/srt/models
+- python/sglang/srt/layers/attention
+
+### AMD optimizations
+[@HaiShaw](https://github.com/HaiShaw) (Henry HAI)
+
+### NPU optimizations
+[@iforgetmyname](https://github.com/iforgetmyname) (Even Zhou)
+
+related files
+- python/sglang/srt/hardware_backend/npu
+
+### CI, Release, Package
+[@Kangyan-Zhou](https://github.com/Kangyan-Zhou) (Kangyan Zhou), [@Fridge003](https://github.com/Fridge003) (Baizhou Zhang)
+
+related files
+- .github/workflows
+
+### Router, API
+[@slin1237](https://github.com/slin1237) (Simo Lin)
+
+related files
+- sgl-model-gateway
+- python/sglang/srt/grpc
+- python/sglang/srt/entrypoints
+
+### Other Notes
+
+Now we have many Merge Oncalls mainly because the CI is flaky and the CODEOWNERS is too coarse-grained.
+In the future, we hope the CI can be improved and we only need bypass rarely. After that, most Merge Oncalls can be converted back to Write and CODEOWNERS.
+
+This list is based on the current situation. If you or someone you know would like to take on more responsibility and are qualified, please ping [Lianmin Zheng](https://github.com/merrymercy) and [Ying Sheng](https://github.com/Ying1123) in the Slack channel. They will start a nomination and internal review process.
+
+## The List of CI Oncalls
+This section lists the oncalls for each hardware platform. The format is @github-username (Slack username).
+
+### NVIDIA GPUs
+[@Kangyan-Zhou](https://github.com/Kangyan-Zhou) (Kangyan Zhou), [@ch-wan](https://github.com/ch-wan) (Cheng Wan), [@HanHan009527](https://github.com/HanHan009527) (hanhan), [@ishandhanani](https://github.com/ishandhanani) (Ishan Dhanani), [@ShangmingCai](https://github.com/ShangmingCai) (Shangming Cai), [@alisonshao](https://github.com/alisonshao) (Alison Shao).
+
+### AMD GPUs
+[@saienduri](https://github.com/saienduri) (Sai Enduri), [@HaiShaw](https://github.com/HaiShaw) (Henry HAI)
+
+### Intel CPU and XPU
+[@mingfeima](https://github.com/mingfeima) (Mingfei Ma), [@DiweiSun](https://github.com/DiweiSun) (Diwei Sun)
+
+### Ascend NPUs
+[@iforgetmyname](https://github.com/iforgetmyname) (Even Zhou)
+
+This list is based on the current situation. If you or someone you know would like to donate machines for CI, they can serve as the CI oncalls for their machines. Please ping [Lianmin Zheng](https://github.com/merrymercy) and [Ying Sheng](https://github.com/Ying1123) in the Slack channel. They will start a nomination and internal review process.
+
+## CI Maintenance Mode
+When the CI is unhealthy (e.g., the scheduled pr-test on `main` is broken for consecutive runs), the project enters **CI Maintenance Mode** by opening [issue #21065](https://github.com/sgl-project/sglang/issues/21065). While active:
+- All PR CI runs are paused. Resources are allocated to PRs that fix the CI.
+- **Merging non-CI-fix PRs is prohibited.** Only PRs that fix the CI may be merged. In severe cases, merge permissions may be revoked.
+
+Maintenance mode ends when `pr-test.yml` is all green on `main` and the issue is closed.
+
+## Suspending Permissions
+If a Merge Oncall bypasses checks to merge a PR that breaks the `main` branch, merges a non-CI-fix PR during CI Maintenance Mode, or repeatedly breaks the CI due to various reasons, their privileges will be suspended for at least two days, depending on the severity of the incident.

third_party/sglang/.github/actions/check-maintenance/action.yml

@@ -0,0 +1,63 @@
+name: Check Maintenance Mode
+description: Blocks CI when maintenance mode is active (issue #21065 is open), unless the PR has the bypass-maintenance label, or env SGLANG_PR_TEST_BYPASS_MAINTENANCE_ON_MAIN=true (PR Test workflow on main only). Merging non-CI-fix PRs is prohibited during maintenance mode; in severe cases, merge permissions may be revoked.
+
+inputs:
+  github-token:
+    description: GitHub token for API access
+    required: false
+    default: ${{ github.token }}
+
+runs:
+  using: composite
+  steps:
+    - name: Check maintenance mode
+      shell: bash
+      env:
+        GH_TOKEN: ${{ inputs.github-token }}
+      run: |
+        MAINTENANCE_ISSUE=21065
+        REPO="${{ github.repository }}"
+        PR_NUMBER="${{ github.event.pull_request.number }}"
+
+        # PR Test workflow only: scheduled runs and runs on main (dispatch / workflow_call) set this env
+        if [[ "${SGLANG_PR_TEST_BYPASS_MAINTENANCE_ON_MAIN:-}" == "true" ]]; then
+          echo "✅ PR Test on main branch; bypassing maintenance gate."
+          exit 0
+        fi
+
+        # Check if maintenance issue is open (fail-open: if API errors, allow CI to proceed)
+        ISSUE_STATE=$(gh issue view "$MAINTENANCE_ISSUE" --repo "$REPO" --json state --jq '.state' 2>/dev/null || echo "UNKNOWN")
+
+        if [[ "$ISSUE_STATE" != "OPEN" ]]; then
+          echo "✅ Maintenance mode is OFF. Proceeding with CI."
+          exit 0
+        fi
+
+        # For PRs, check if bypass-maintenance label is present
+        if [[ -n "$PR_NUMBER" ]]; then
+          HAS_BYPASS=$(gh pr view "$PR_NUMBER" --repo "$REPO" --json labels --jq '[.labels[].name] | map(select(. == "bypass-maintenance")) | length' 2>/dev/null || echo "0")
+          if [[ "$HAS_BYPASS" -gt 0 ]]; then
+            echo "✅ PR #$PR_NUMBER has 'bypass-maintenance' label. Bypassing maintenance mode."
+            exit 0
+          fi
+        fi
+
+        MSG=$(printf "%s\n" \
+          "## ⚠️ CI Maintenance Mode is Active" \
+          "The CI infrastructure is currently under maintenance." \
+          "All PR CI runs are paused until maintenance is complete." \
+          "**Merging non-CI-fix PRs is prohibited during maintenance mode.** In severe cases, merge permissions may be revoked." \
+          "You might also experience unexpected failures during this period." \
+          "The team is working on the issue and will update the status as soon as possible." \
+          "" \
+          "What should you do?" \
+          "- **Do NOT merge non-CI-fix PRs** until maintenance mode is lifted" \
+          "- Check back later (~12 hours)" \
+          "- Follow CI Maintenance Mode issue: https://github.com/$REPO/issues/$MAINTENANCE_ISSUE for status updates")
+
+        echo "$MSG" >> "$GITHUB_STEP_SUMMARY"
+        while IFS= read -r line; do
+          echo "::error::$line"
+        done <<< "$MSG"
+
+        exit 1

third_party/sglang/.github/actions/check-stage-health/action.yml

@@ -0,0 +1,50 @@
+name: Check Stage Health
+description: Fail fast if any job in the current workflow run has already failed. Auto-skips for scheduled runs.
+
+inputs:
+  github-token:
+    description: 'GitHub token for API calls'
+    required: false
+    default: ${{ github.token }}
+
+runs:
+  using: composite
+  steps:
+    - name: Check stage health
+      uses: actions/github-script@v7
+      env:
+        SKIP_STAGE_HEALTH_CHECK: ${{ env.SKIP_STAGE_HEALTH_CHECK }}
+      with:
+        github-token: ${{ inputs.github-token }}
+        script: |
+          // Skip when explicitly requested via env var (e.g. release branch cut)
+          if (process.env.SKIP_STAGE_HEALTH_CHECK === 'true') {
+            core.info('Skipping health check (SKIP_STAGE_HEALTH_CHECK=true)');
+            return;
+          }
+
+          // Skip for scheduled runs — they should collect all failures, not fast-fail
+          if (context.eventName === 'schedule') {
+            core.info('Skipping health check for scheduled run');
+            return;
+          }
+
+          const jobs = await github.paginate(github.rest.actions.listJobsForWorkflowRun, {
+            owner: context.repo.owner,
+            repo: context.repo.repo,
+            run_id: context.runId,
+            per_page: 100,
+          });
+          // Find jobs that failed from a real error, not from fast-fail cascade
+          const rootCauseFailures = jobs.filter(j => {
+            if (j.status !== 'completed' || j.conclusion !== 'failure') return false;
+            // If the failing step is the health check, it's a cascade — skip it
+            const failedStep = (j.steps || []).find(s => s.conclusion === 'failure');
+            if (failedStep && (failedStep.name.includes('check-stage-health') || failedStep.name.includes('Check stage health'))) {
+              return false;
+            }
+            return true;
+          });
+          if (rootCauseFailures.length > 0) {
+            core.setFailed(`Fast-fail: skipping — root cause job(s): ${rootCauseFailures.map(j => j.name).join(', ')}`);
+          }

third_party/sglang/.github/actions/upload-cuda-coredumps/action.yml

@@ -0,0 +1,27 @@
+name: Upload CUDA Coredumps
+description: Upload CUDA coredump files as artifacts and clean up the directory.
+
+inputs:
+  artifact-suffix:
+    description: Suffix appended to the artifact name (e.g. matrix partition id)
+    required: false
+    default: ""
+  retention-days:
+    description: Number of days to retain the artifact
+    required: false
+    default: "7"
+
+runs:
+  using: composite
+  steps:
+    - name: Upload CUDA coredumps
+      uses: actions/upload-artifact@v4
+      with:
+        name: cuda-coredumps-${{ github.job }}${{ inputs.artifact-suffix && format('-{0}', inputs.artifact-suffix) }}
+        path: ${{ env.SGLANG_CUDA_COREDUMP_DIR || '/tmp/sglang_cuda_coredumps' }}/
+        retention-days: ${{ inputs.retention-days }}
+        if-no-files-found: ignore
+
+    - name: Cleanup CUDA coredumps
+      shell: bash
+      run: rm -rf "${{ env.SGLANG_CUDA_COREDUMP_DIR || '/tmp/sglang_cuda_coredumps' }}"