docs(kvc): record real Ali KVC experiment results

feat(kvc): add real Ali replay workflow
docs(kvc): agentic-fit analysis, refactor plan, validation report
2026-05-12 05:28:06 +00:00 · 2026-05-12 05:28:00 +00:00 · 2026-05-06 21:30:11 +08:00 · 2026-05-06 21:29:56 +08:00 · 2026-05-06 21:29:46 +08:00 · 2026-05-06 21:29:30 +08:00
4369 changed files with 1219433 additions and 2 deletions
--- a/.gitignore
+++ b/.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
--- a/AGENTS.md
+++ b/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.
--- a/README.md
+++ b/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__`、虚拟环境。
--- a/docs/AGENTIC_FIT_ANALYSIS_ZH.md
+++ b/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 实验
--- a/docs/KVCACHE_CENTRIC_PROGRESS_ZH.md
+++ b/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。
--- a/docs/KVC_DEBUG_JOURNEY_V1_TO_V5.md
+++ b/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 决策是否在说谎。
--- a/docs/PROJECT_OVERVIEW.md
+++ b/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 预测。
--- a/docs/REAL_ALI_KVC_EXPERIMENT_LOG_ZH.md
+++ b/docs/REAL_ALI_KVC_EXPERIMENT_LOG_ZH.md
@@ -0,0 +1,514 @@
+# Real Ali KVC 实验日志
+
+**分支**：`kvc-real-ali-iter-v1`，从 `kvc-debug-journey-v1-to-v4` checkout 出来。
+**日期**：2026-05-11/12。
+**环境**：单机 8x NVIDIA H20，SGLang xPyD，模型 `/home/admin/cpfs/wjh/models/Qwen/Qwen3-Coder-30B-A3B-Instruct`。
+**真实 trace**：`/home/admin/cpfs/wjh/ali-trace/trace-qwen3-coder-formatted/041715-041717.jsonl`。
+
+本日志记录真实 Ali workload 上的 KVC pd-hybrid 迭代。结论只按当前证据成立；`time-scale=10` smoke 和 KVC-friendly slice 不作为 full workload headline。
+
+## 1. 当前最新进展
+
+已新增真实 Ali trace 的固定样本和 sweep 管线：
+
+- `scripts/prepare_real_ali_samples.py`：从真实 Ali trace 生成可复现实验样本，保留真实 input/output/hash_ids/timestamp，可选择 rebase timestamp。
+- `scripts/sweep_real_ali_kvc.sh`：对同一 prebuilt sample 依次跑 DP cache-aware、PD-disaggregation、KVC、KVC+backpressure。
+- `benchmark-live --use-trace-as-sample`：直接 replay 指定 trace，避免不同策略重新采样导致不可比。
+- `replay-progress.jsonl` heartbeat：后续长跑会每 30s 写客户端侧进度，不轮询 `/server_info`，避免扰动 scheduler。
+- `prepare_real_ali_samples.py --max-sampled-duration-s`：为快速 smoke 生成 capped sample；只用于迭代，不用于 headline。
+
+已经完成的真实 Ali KVC-fit smoke：
+
+- 样本：`outputs/real-ali-kvc-iter/samples-balanced/ali-kvc-fit-smallappend.jsonl`
+- 179 requests，64 sessions，全部 multi-turn；turn2+ 共 115 个，direct-eligible ratio 100%。
+- `time-scale=10`，concurrency 32。
+- DP cache-aware、PD-disaggregation、KVC no-backpressure、KVC+backpressure 均已完成。
+
+## 2. 全量 Ali trace 画像
+
+`outputs/real-ali-kvc-iter/ali-full-profile.json` 显示：
+
+| 指标 | 数值 |
+|---|---:|
+| requests | 763,727 |
+| sessions | 555,905 |
+| multi-turn sessions | 39,247 |
+| turn2+ requests | 207,822 |
+| turn2+ direct-eligible ratio | 82.95% |
+| input p50 / p90 / p99 | 4,329 / 51,067 / 112,955 tokens |
+| output p50 / p90 / p99 | 93 / 826 / 5,616 tokens |
+| append p50 / p90 / p99 | 303 / 2,879 / 17,885 tokens |
+| inter-turn gap p50 / p90 / p99 | 4.65s / 38.68s / 1,133s |
+
+这个 profile 说明 KVC 有真实适用面：turn2+ 的 hash overlap 和小 append 很常见。但 full workload 里 single-turn session 极多，KVC 收益会被显著稀释；因此必须分 slice 报告，不能只报 KVC-fit 子集。
+
+## 3. 已跑样本
+
+### Continuous 15min cold-window session sample
+
+路径：`outputs/real-ali-kvc-iter/samples-window-900s-600req/ali-window.jsonl`
+
+- 600 requests，439 sessions，32 multi-turn sessions。
+- rebased duration：886.544s，覆盖约 15min。
+- turn2+ requests：161，direct-eligible：143，ratio 88.8%。
+- input p50 / p90 / p99：3,871 / 68,234 / 98,131。
+- output p50 / p90 / p99：85 / 712 / 5,195。
+- append p50 / p90 / p99：274 / 2,202 / 16,120。
+- inter-turn gap p50 / p90 / p99：4.656s / 19.376s / 63.575s。
+
+这是对 179-request KVC-fit smoke 的替代验证样本。它按 900s 窗口分成 15 个时间桶，轮转选择窗口内从 root 开始的整 session，直到达到 600 requests。这样避免 parent 缺失导致 `load_trace()` 把真实 session 切碎，也让请求覆盖整个 15min，而不是只取窗口开头 600 条。
+
+重要边界：它是 **cold-window / new-session-only** sample，不是完整 raw production window；它排除了窗口开始前已经活跃的 ongoing sessions。因此可以用于“600+ 请求、15min、真实混合负载”的稳定性验证，但不能单独代表全量 Ali production window。
+
+### KVC-fit small append
+
+路径：`outputs/real-ali-kvc-iter/samples-balanced/ali-kvc-fit-smallappend.jsonl`
+
+- 179 requests，64 sessions。
+- input p50 / p90：6,446 / 15,491。
+- output p50 / p90：112 / 1,159。
+- append p50 / p90：215 / 855。
+- overlap ratio p50 / p90：0.875 / 0.938。
+
+这是 KVC-friendly slice，用来验证机制上限和 microbenchmark 是否能迁移到真实 token/hash 序列。
+
+### Representative-mt / early multi-turn balanced
+
+路径：`outputs/real-ali-kvc-iter/samples-balanced/ali-representative-mt.jsonl`
+
+- 460 requests，64 sessions。
+- input p50 / p90：41,175 / 98,621。
+- append p50 / p90 / p99：272 / 1,979 / 13,900。
+
+这个样本更接近真实 multi-turn 压力，后续用于验证大上下文、大 resident KV 下是否仍能稳定。但它当前实现是“从 start_time 后取最早 64 个 multi-turn session”，不是严格随机或分层 representative；正式 headline 需要按 input/append/output/gap 分层抽样。
+
+### Capped smoke samples
+
+为避免少数真实长 gap 让 smoke 浪费大量 wall time，新增：
+
+- `outputs/real-ali-kvc-iter/samples-balanced-cap120s/ali-kvc-fit-smallappend.jsonl`：177 requests，64 sessions，duration 65.859s。
+- `outputs/real-ali-kvc-iter/samples-balanced-cap120s/ali-representative-mt.jsonl`：359 requests，64 sessions，duration 117.366s。
+
+这些样本去掉了 KVC-fit 原样本末尾 timestamp 3613s 和 5414s 的两个请求，因此只能用于快速工程迭代；正式对比仍应使用完整样本或真实连续窗口。
+
+## 4. 当前结果
+
+### 4.1 DP cache-aware vs KVC+backpressure, KVC-fit, time-scale=10
+
+| 策略 | Requests | Errors | Trunc | E2E mean | E2E p50 | E2E p90 | E2E p99 | TTFT mean | TTFT p50 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---:|
+| 8-way DP cache-aware | 179 | 0 | 0 | 6.603s | 3.126s | 17.639s | 34.582s | 1.112s | 1.052s |
+| KVC 2P6D + worker admission + backpressure | 179 | 0 | 0 | 4.443s | 2.076s | 13.288s | 21.202s | 0.700s | 0.154s |
+
+Paired comparison（KVC - DP）：
+
+- overall E2E mean delta：-2.161s；p50 delta：-1.427s；152/179 wins。
+- turn2+ direct 子集：mean delta -2.503s；p50 delta -1.508s；103/115 wins。
+- turn2+ TTFT mean delta：-0.930s；p50 delta -0.887s。
+
+执行路径：
+
+- KVC turn1 seed：64 requests。
+- `kvcache-direct-to-d-session`：115 requests。
+- session reused：115。
+- actual KV transfer blocks：623。
+
+结构日志：
+
+- admission probes：179，全为 `ok`。
+- transfer queue depth：p50=0，p90=2，max=3。
+- backpressure event：0。
+
+解释：这轮证明的是 **KVC direct-to-D/session reuse** 在真实 Ali KVC-fit slice 上有正信号；不是证明 backpressure 有效，因为没有触发 backpressure。
+
+### 4.2 PD-disaggregation baseline, KVC-fit, time-scale=10
+
+| 策略 | Requests | Errors | Trunc | E2E mean | E2E p50 | E2E p90 | E2E p99 | TTFT mean | TTFT p50 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---:|
+| PD-disaggregation 2P6D | 179 | 0 | 0 | 7.850s | 6.306s | 15.192s | 22.405s | 4.994s | 5.336s |
+
+Paired comparison（PD - DP）：
+
+- overall E2E mean delta：+1.247s。
+- p50 delta：+2.231s。
+- 46/179 faster，133/179 slower。
+
+解释：在这个 KVC-fit slice 上，普通 PD-disaggregation 明显弱于 8-way DP cache-aware。它付出了 P->D transfer 和拆分调度成本，却没有 KVC direct-to-D 的 bypass 收益。
+
+### 4.3 KVC no-backpressure 消融, KVC-fit, time-scale=10
+
+| 策略 | Requests | Errors | Trunc | E2E mean | E2E p50 | E2E p90 | E2E p99 | TTFT mean | TTFT p50 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---:|
+| KVC 2P6D worker admission, no backpressure | 179 | 0 | 0 | 4.404s | 1.936s | 13.200s | 21.326s | 0.604s | 0.139s |
+
+Paired comparison：
+
+- KVC no-BP vs DP：mean delta -2.200s，p50 delta -1.434s，153/179 wins。
+- KVC no-BP vs PD-disaggregation：mean delta -3.447s，p50 delta -3.514s，163/179 wins。
+- KVC no-BP vs KVC+BP：mean delta -0.039s，p50 delta -0.005s，92/179 wins。
+
+结构分析：
+
+- direct-to-D rate：64.25%。
+- admission probes：179，全为 `ok`。
+- transfer queue depth：p50=0，p90=2，max=3。
+- pause_ms 全 0，backpressure event 0。
+
+解释：no-backpressure 与 +backpressure 几乎等价，说明本 slice 没有 D 压力；本轮提升来自 direct-to-D，不来自反压。
+
+### 4.4 Continuous 15min / 600-request window, time-scale=1
+
+样本：`outputs/real-ali-kvc-iter/samples-window-900s-600req/ali-window.jsonl`
+
+重要边界：这是 cold-window / new-session-only session sample，不是完整 raw window。它覆盖约 15min，且 `missing_parent_count=0`，但排除了窗口开始前已活跃的 ongoing sessions。
+
+运行结果：
+
+| 策略 | Requests | Errors | Trunc | E2E mean | E2E p50 | E2E p90 | E2E p99 | TTFT mean | TTFT p50 | TTFT p90 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---:|---:|
+| DP cache-aware 8-way | 600 | 1 | 0 | 13.942s | 5.222s | 29.299s | 151.183s | 6.162s | 1.746s | 19.176s |
+| PD-disaggregation 2P6D | 600 | 1 | 0 | 40.886s | 40.018s | 84.681s | 113.460s | 38.545s | 37.782s | 81.852s |
+| KVC 2P6D mem_fraction_static=0.82 | 600 | 53 | 0 | 12.386s | 4.225s | 37.998s | 78.234s | 10.078s | 2.674s | 27.774s |
+
+KVC 默认启动失败：
+
+- 默认 KVC 2P6D 在 H20 上两次启动 OOM，均未进入 replay。
+- 日志显示 decode/prefill worker 启动时只剩约 526MB，模型加载阶段 OOM。
+- `--load-format layered` 不支持 Qwen3-Coder-30B-A3B。
+- 使用 `--mem-fraction-static 0.82` 后 KVC 能启动并完成 replay，但这降低了 KV pool 容量，因此这轮 KVC 是 memory-constrained rerun。
+- 尝试 `KVC_SEED_MIN_TURN_ID=2` + `mem_fraction_static=0.82` 时，启动阶段 scheduler 被 SIGKILL，疑似 OS OOM killer，未进入 replay。
+
+Paired comparison（只在两边都有 latency 的 547 个 paired request 上计算）：
+
+- KVC vs DP：mean delta -1.335s，p50 delta -0.055s，p90 delta +19.371s，284 wins / 263 losses。
+- KVC vs PD：mean delta -28.341s，p50 delta -25.687s，p90 delta +2.834s，465 wins / 82 losses。
+
+KVC 结构数据：
+
+- execution modes：388 `pd-router-turn1-seed`，90 `kvcache-direct-to-d-session`，67 `pd-router-fallback-large-append-session-cap`，1 `pd-router-large-append-reseed`，1 `pd-router-turn1-d-backpressure`，53 `kvcache-centric` error rows。
+- direct-to-D rate：15.0%。
+- direct-to-D session 分布：413/439 sessions 在 0-20% direct rate；只有 6 sessions 在 80-100%。
+- admission probes：533；reason `ok` 531，`no-space` 2；queue depth p50=0，p90=2，max=5。
+- pause hint 非零 20 次，但没有 backpressure event，因为本轮 no-BP。
+
+KVC error breakdown：
+
+- 50 `ReadTimeout`。
+- 2 `HTTPStatusError 400 Bad Request` on `open_session`。
+- 1 context length error：同 DP/PD 的 `input_length=310521 > 262144`。
+- 错误主要集中在 turn1：50 turn1，3 turn2+。
+
+解释：
+
+1. KVC 相对普通 PD 仍明显更好，说明普通 P->D disaggregation 在真实 600-request 窗口上成本很高。
+2. KVC 相对 DP 只在 clean request 的 mean/p50 上有小幅正信号，但 p90 变差，而且 error_count 从 DP 的 1 增到 53。
+3. 因此在这个 600-request / 15min window 上，**KVC 不能算稳定提升系统**。主要问题不是 direct-to-D 快路径无效，而是该快路径覆盖率只有 15%，并且 turn1 seed / session admission / memory-constrained KV pool 引入大量 timeout。
+4. 这直接修正 179-request KVC-fit smoke 的结论：小样本证明 KVC 适用 slice 存在；600-request mixed window 证明当前实现还不能稳定服务真实混合 workload。
+
+## 5. 是否已经相对 pd-colocation/pd-disaggregation 取得提升
+
+当前只能下这个限定结论：
+
+1. **相对 PD-disaggregation：已经取得清晰提升。**
+   PD-disaggregation p50 6.306s，KVC no-BP p50 1.936s，KVC+BP p50 2.076s；TTFT p50 5.336s vs 0.139s/0.154s。收益主要来自 turn2+ 直接打到已有 D session，避免每轮 P 全量 prefill 和 P->D KV transfer。
+
+2. **相对强 DP cache-aware：在 KVC-fit slice 上有提升。**
+   KVC no-BP 和 KVC+BP overall mean/p50/p90/p99 都优于 DP，并且 paired wins 分别是 153/179 和 152/179。但这是 KVC-friendly、全 multi-turn、turn2+ 100% direct-eligible 的 slice，不代表 full Ali workload。
+
+3. **相对 full workload：尚未证明。**
+   全量 Ali 里 single-turn 占多数，且长上下文和长尾 output 较多。KVC 的收益面会被 single-turn 稀释，D resident KV 容量和 tail 稳定性会成为更强约束。
+
+4. **相对 600-request / 15min mixed window：尚未取得稳定提升。**
+   KVC clean E2E mean/p50 有正信号，但 error_count=53/600，p90 paired delta 相对 DP 变差。按“E2E + error/truncation”标准，这不能算系统性胜出。
+
+## 6. 提升来自哪里
+
+主要收益链路：
+
+1. turn1 seed 在 D 上建立 session。
+2. turn2+ 若 append 小、hash overlap 高，直接走 `kvcache-direct-to-d-session`。
+3. direct-to-D 避免 P worker 参与，不走 P->D KV transfer。
+4. D 只对 append suffix 做少量 prefill，已有前缀 KV 直接复用。
+
+这带来两个可观测收益：
+
+- TTFT 大幅下降：turn2+ direct 子集 TTFT mean 从 DP 的约 1.04s 降到约 0.112s。
+- E2E 下降：direct 子集 mean E2E 降低约 2.50s。
+
+另外，KVC 的 cached_tokens 统计显著更高：KVC mean cached tokens 5,992，DP mean 228。这说明它确实复用了大段真实前缀 KV。
+
+## 7. 遇到的问题与修复
+
+### 问题 1：通用 sampler 会被单个长 session 主导
+
+现象：真实 Ali session 分布长尾明显，duration-oriented 采样容易选出不均衡样本，导致策略比较不可重复或不代表多 session 竞争。
+
+修复：新增 `scripts/prepare_real_ali_samples.py`，按 session 上限和每 session turn 上限生成 balanced sample，并保留真实 token/hash/timestamp。
+
+### 问题 2：不同策略重新采样导致不可比
+
+现象：`benchmark-live` 原本会按参数重新采样，不同策略可能 replay 不同请求。
+
+修复：新增 `--use-trace-as-sample`，所有策略 copy 并 replay 同一个 prebuilt sample；后续 paired comparison 才有意义。
+
+### 问题 3：长 trace replay 中途没有进度
+
+现象：`request-metrics.jsonl` 和 summary 只在 replay 结束后写出，跑真实 pacing 时很难判断是正常等待还是卡住。
+
+修复：新增 `replay-progress.jsonl` heartbeat，每 30s 写 submitted/completed/inflight/errors/execution_modes。它只使用客户端本地状态，不访问 `/server_info`。
+
+### 问题 4：`/server_info` polling 会扰动 scheduler
+
+现象：旧 profiling 里 1Hz polling 曾明显改变错误数。真实 performance run 如果持续 poll pool，会把测量工具变成干扰源。
+
+修复：`scripts/sweep_real_ali_kvc.sh` 默认关闭 pool polling。容量类问题依赖结构日志和必要时单独 profile run，不混入 headline performance run。
+
+### 问题 5：backpressure smoke 没有触发 backpressure
+
+现象：KVC-fit smoke 中 transfer queue max 只有 3，所有 admission reason 都是 `ok`，pause_ms 全 0。
+
+结论：这轮不能证明 backpressure 有效，只能证明 direct-to-D 有效。需要更高 session 数、更大 resident KV 或更强并发的压力样本专门验证 backpressure。
+
+### 问题 6：环境和旧报告不一致
+
+现象：旧文档写的是 H100，本轮真实环境是 H20；模型路径也在 `/home/admin/cpfs/wjh/models/...`。
+
+处理：本日志按 H20 记录；跨文档比较时只看机制趋势，不把 H100/H20 的绝对 latency 混为同一实验。
+
+### 问题 7：continuous window 可能截断 session ancestry
+
+现象：按 timestamp 直接截窗口可能留下 parent turn 在窗口外的请求。对 KVC 来说，这会让 session reuse/turn chain 与真实 workload 不一致。
+
+处理：当前 continuous window 只作为待改进候选，不作为正式 headline。正式窗口需要保留 warmup ancestors，或显式保留原始 session chain 信息。
+
+## 8. 如果后续 full workload 效果不好，当前假设
+
+可能不是实现小 bug，而是方案适用面和资源约束共同导致：
+
+1. **single-turn 稀释收益**：全量 Ali session 中 single-turn 占多数，KVC seed 只带来成本，没有 turn2+ reuse。
+2. **长上下文挤占 D KV 池**：input p90 51K、p99 113K，resident KV 长尾会限制 D 上可同时保留的 session。
+3. **direct 不是免费 lunch**：turn1 seed、admission probe、session lifecycle 都有额外成本；只有后续 turns 充分复用时才摊薄。
+4. **D 端容量和 eviction 仍是核心风险**：旧 SWE 实验已经显示 session pinning + D 容量盲选会造成 starvation；early multi-turn balanced 样本可能复现。
+5. **普通 PD-disaggregation 很弱**：如果 KVC fallback 频繁退回普通 PD 路径，整体会被 P->D transfer 和高 TTFT 拖垮。
+6. **H20 显存余量不足会改变 KVC 条件**：默认 KVC 2P6D 启动 OOM，必须降 `mem_fraction_static` 才能完成 600-request run；这会进一步降低 D KV pool，放大 session-cap 和 timeout。
+
+## 9. 下一步验证顺序
+
+1. 补 sticky/session-affinity baseline，拆出“粘到同一个 D”和“KVC direct bypass”的贡献。
+2. 补 KVC `seed-min-turn-id=2` 或 no-turn1-seed，验证 turn1 seed 成本是否值得。
+3. 在 early multi-turn balanced 样本上跑 DP / PD / KVC no-BP / KVC+BP，验证大上下文真实 multi-turn 压力。
+4. 选小固定样本跑 `time-scale=1`，避免只在压缩 replay 条件下成立。
+5. 做包含 single-turn 的 continuous window，并处理窗口内 parent turn 缺失问题，再按 full Ali 分布加权报告。
+6. 对最终候选配置做 N>=3 rerun，报告方差；N=1 只作为 smoke。
+7. 针对 600-request window 优先跑 `seed-min-turn-id=2`，减少 single-turn turn1 seed；目标是先把 53/600 errors 降到接近 DP 的 1/600，再讨论 latency。
+   - 当前第一次尝试未进入 replay，启动阶段疑似 OS OOM；需要先解决 H20 启动显存/系统内存稳定性，或者降低 worker 数/模型内存占用。
+
+## 10. KVC error 根因与 multi-turn-only 验证准备
+
+用户指出 179-request run 不够，并要求至少 15min / 600+ 请求；当前正式问题定位基于
+`outputs/real-ali-kvc-iter/runs/window900s-600req-ts1-kvc-mem082/kvcache-centric-kv-aware-worker-admission-20260511T093601Z`。
+
+### 10.1 为什么 KVC 有大量 error
+
+该 run 为 600 requests，KVC mem0.82 有 53 errors：
+
+- 50 个 `ReadTimeout`。
+- 2 个 `/open_session` HTTP 400。
+- 1 个真实超上下文错误：input 310,521 > model context 262,144。
+
+按 turn 看，50/53 errors 在 turn1。按 structural admission 看，绝大多数失败请求在
+`structural/admission-events.jsonl` 中已经被 D 端 admission 判定 `can_admit=true`，所以这不是单纯的
+`d-session-cap` 或 `no-space`。主要失败点是 turn1 seed 进入 KVC seeded path 后，在
+P/D streaming session bootstrap、P->D transfer 或 router streaming 过程中超时；而混合真实窗口中 single-turn session 很多，
+这些 turn1 seed 对大多数 session 没有后续复用收益。
+
+结论：当前 KVC error 的主因是 **对 single-turn / 未知是否多轮的 session 做了过多 turn1 seed**，它把大量新 session 推进
+KVC control-plane 和 seeded router 路径，增加超时和 session lifecycle 残留；不是 direct-to-D fast path 本身出错。
+
+### 10.2 已做修复/消融开关
+
+代码与脚本修复：
+
+- `scripts/sweep_real_ali_kvc.sh` 新增 `KVC_SEED_ONLY_MULTITURN=1`，会传入
+  `--kvcache-seed-only-multiturn-sessions`。这是 oracle 消融，用来验证“只 seed 会有后续 turn 的 session”能否消除 turn1 seed 错误。
+- `src/agentic_pd_hybrid/replay.py` 对 `/open_session` 400 增加 close+retry 一次，并写
+  `structural/session-lifecycle.jsonl`。这是 lifecycle 健壮性修复，目标是处理 timeout 后服务端残留 session 导致的
+  “already exists” 400，不改变 routing policy。
+- `scripts/prepare_real_ali_samples.py` 新增 `--window-min-turns` 和 `--window-output-name`，用于生成可复现的 multi-turn-only window 样本。
+
+验证：
+
+- `uv run python -m py_compile scripts/prepare_real_ali_samples.py src/agentic_pd_hybrid/replay.py src/agentic_pd_hybrid/benchmark.py src/agentic_pd_hybrid/cli.py`
+- `bash -n scripts/sweep_real_ali_kvc.sh`
+
+### 10.3 已生成 multi-turn-only 样本
+
+样本路径：
+
+`outputs/real-ali-kvc-iter/samples-window-900s-600req-multiturn/ali-window-multiturn.jsonl`
+
+生成命令：
+
+```bash
+uv run python scripts/prepare_real_ali_samples.py \
+  --trace /home/admin/cpfs/wjh/ali-trace/trace-qwen3-coder-formatted/041715-041717.jsonl \
+  --output-root outputs/real-ali-kvc-iter/samples-window-900s-600req-multiturn \
+  --window-duration-s 900 \
+  --window-target-requests 600 \
+  --window-buckets 15 \
+  --window-min-turns 2 \
+  --window-output-name ali-window-multiturn.jsonl \
+  --profiles representative-mt \
+  --max-sessions 64 \
+  --max-turns-per-session 12
+```
+
+样本 profile：
+
+- 626 requests，107 sessions，107 个都是 multi-turn sessions。
+- sampled duration 889.341s。
+- turn2+ = 519。
+- direct-eligible turn2+ = 473 / 519 = 91.1%。
+- missing parent = 0。
+- input p50/p90/p99 = 26,846 / 91,596 / 123,898 tokens。
+
+这个 case 是“过滤掉 single-turn 的多轮压力切片”，不能替代 full mixed workload，但可以回答：
+如果 workload 确实以多轮 coding agent session 为主，KVC 的 direct-to-D 覆盖率和稳定性是否接近 microbenchmark。
+
+### 10.4 GPU 资源阻塞
+
+截至本次记录，8 张 GPU 均被另一组 `vllm serve` 进程占用，每张约 82GB / 98GB，端口为 51000-51007。
+这些不是本 repo 的 SGLang/benchmark 进程，因此未启动新的性能 run，避免把资源冲突误判为 KVC 策略失败。
+
+GPU 释放后，优先跑两组：
+
+```bash
+# 混合真实窗口：验证 seed-only-multiturn 是否把 53/600 errors 降下来
+TRACE=outputs/real-ali-kvc-iter/samples-window-900s-600req/ali-window.jsonl \
+OUT_ROOT=outputs/real-ali-kvc-iter/runs/window900s-600req-ts1-kvc-seedonly-mt-mem082 \
+RUNS="kvc" \
+TIME_SCALE=1 \
+CONCURRENCY=32 \
+REQUEST_TIMEOUT_S=600 \
+STACK_TIMEOUT_S=1800 \
+EXTRA_SERVER_ARGS="--mem-fraction-static 0.82" \
+KVC_SEED_ONLY_MULTITURN=1 \
+bash scripts/sweep_real_ali_kvc.sh
+
+# 多轮-only workload：DP vs KVC，对照过滤 workload 是否能复现 microbenchmark 收益
+TRACE=outputs/real-ali-kvc-iter/samples-window-900s-600req-multiturn/ali-window-multiturn.jsonl \
+OUT_ROOT=outputs/real-ali-kvc-iter/runs/window900s-600req-multiturn-ts1-mem082 \
+RUNS="dp kvc" \
+TIME_SCALE=1 \
+CONCURRENCY=32 \
+REQUEST_TIMEOUT_S=600 \
+STACK_TIMEOUT_S=1800 \
+EXTRA_SERVER_ARGS="--mem-fraction-static 0.82" \
+KVC_SEED_ONLY_MULTITURN=1 \
+bash scripts/sweep_real_ali_kvc.sh
+```
+
+### 10.5 multi-turn-only 启动尝试被 GPU 占用阻塞
+
+用户要求启动 multi-turn-only 的 `pd-disaggregation` vs `kvcache-centric` 对比。启动前检查发现 8 张 GPU 均被外部
+`vllm serve` 进程占用，每张约 84GB / 98GB，端口为 51000-51007。该进程不属于本 repo 的 SGLang/benchmark run。
+
+因此本次没有强行启动 SGLang。原因是剩余显存不足以启动 2P6D 或 8-worker 对照，强行运行只会得到初始化 OOM 或不稳定超时，
+不能用于判断 KVC pd-hybrid 是否优于 pd-disaggregation。
+
+资源释放后要运行的 multi-turn-only 对比命令：
+
+```bash
+TRACE=outputs/real-ali-kvc-iter/samples-window-900s-600req-multiturn/ali-window-multiturn.jsonl \
+OUT_ROOT=outputs/real-ali-kvc-iter/runs/window900s-600req-multiturn-ts1-pd-vs-kvc-mem082 \
+RUNS="pd kvc" \
+TIME_SCALE=1 \
+CONCURRENCY=32 \
+REQUEST_TIMEOUT_S=600 \
+STACK_TIMEOUT_S=1800 \
+EXTRA_SERVER_ARGS="--mem-fraction-static 0.82" \
+KVC_SEED_ONLY_MULTITURN=1 \
+bash scripts/sweep_real_ali_kvc.sh
+```
+
+### 10.6 multi-turn-only PD vs KVC 正式结果
+
+资源释放后已启动并完成 multi-turn-only 对比。运行命令：
+
+```bash
+TRACE=outputs/real-ali-kvc-iter/samples-window-900s-600req-multiturn/ali-window-multiturn.jsonl \
+OUT_ROOT=outputs/real-ali-kvc-iter/runs/window900s-600req-multiturn-ts1-pd-vs-kvc-mem082 \
+RUNS="pd kvc" \
+TIME_SCALE=1 \
+CONCURRENCY=32 \
+REQUEST_TIMEOUT_S=600 \
+STACK_TIMEOUT_S=1800 \
+EXTRA_SERVER_ARGS="--mem-fraction-static 0.82" \
+KVC_SEED_ONLY_MULTITURN=1 \
+bash scripts/sweep_real_ali_kvc.sh
+```
+
+Run 目录：
+
+- PD：`outputs/real-ali-kvc-iter/runs/window900s-600req-multiturn-ts1-pd-vs-kvc-mem082/pd-disaggregation-kv-aware-20260512T030433Z`
+- KVC：`outputs/real-ali-kvc-iter/runs/window900s-600req-multiturn-ts1-pd-vs-kvc-mem082/kvcache-centric-kv-aware-worker-admission-20260512T040444Z`
+
+样本仍是 626 requests、107 sessions、889.341s，全部为 multi-turn session。
+
+| 策略 | Requests | Errors | Trunc | E2E mean | E2E p50 | E2E p90 | E2E p99 | TTFT mean | TTFT p50 | TTFT p90 |
+|---|---:|---:|---:|---:|---:|---:|---:|---:|---:|---:|
+| PD-disaggregation 2P6D | 626 | 0 | 0 | 97.013s | 70.243s | 214.309s | 308.406s | 94.506s | 69.048s | 212.528s |
+| KVC 2P6D worker admission, no BP, seed-only-multiturn | 626 | 39 | 0 | 43.362s | 8.239s | 135.289s | 236.475s | 40.578s | 1.442s | 132.233s |
+
+Paired comparison 只在 KVC 成功且 PD 也有 latency 的 587 个 request 上计算：
+
+- PD same-request E2E mean/p50/p90/p99：97.457s / 70.514s / 214.095s / 309.362s。
+- KVC same-request E2E mean/p50/p90/p99：43.362s / 8.239s / 135.930s / 237.283s。
+- mean E2E reduction：55.5%。
+- absolute mean improvement：54.095s。
+- wins/losses：472 / 115。
+
+按 KVC execution mode 拆分：
+
+| KVC mode | Count | KVC mean | PD same mean | Reduction |
+|---|---:|---:|---:|---:|
+| `kvcache-direct-to-d-session` | 286 | 2.255s | 92.944s | 97.6% |
+| `pd-router-fallback-large-append-session-cap` | 169 | 88.869s | 113.614s | 21.8% |
+| `pd-router-d-session-reseed` | 25 | 143.456s | 106.501s | -34.7% |
+| `pd-router-large-append-reseed` | 19 | 47.631s | 88.981s | 46.5% |
+| `pd-router-turn1-seed` | 78 | 55.974s | 73.050s | 23.4% |
+
+按 turn 深度拆分：
+
+- turn2+：504 successful paired requests，KVC mean 40.791s vs PD mean 101.055s，reduction 59.6%。
+- turn>=5：299 successful paired requests，KVC mean 34.121s vs PD mean 104.697s，reduction 67.4%。
+- turn>=10：161 successful paired requests，KVC mean 39.027s vs PD mean 86.548s，reduction 54.9%。
+
+KVC execution modes：
+
+- `kvcache-direct-to-d-session`：286。
+- `pd-router-fallback-large-append-session-cap`：169。
+- `pd-router-turn1-seed`：78。
+- `pd-router-d-session-reseed`：25。
+- `pd-router-large-append-reseed`：19。
+- `pd-router-fallback-no-d-capacity`：4。
+- `pd-router-turn1-d-backpressure`：5。
+- `pd-router-d-session-reseed-after-eviction`：1。
+- error rows：39，记录为 `kvcache-centric`。
+
+KVC 的收益来源非常清楚：286 个 direct-to-D request 的 same-request mean 从 PD 的 92.944s 降到 2.255s，基本复现了 microbenchmark 的核心机制收益。它跳过 P worker 和 P->D KV transfer，只在已有 D session 上处理 append suffix。总体 actual KV transfer blocks 从 PD same-success 的 4436 降到 KVC success 的 3827；summary 口径下 KVC total actual KV transfer blocks 为 3827，低于 PD 的 5276。
+
+但这轮仍不能作为“稳定生产级胜出”结论：
+
+1. KVC 仍有 39/626 errors，error rate 6.23%，PD 为 0。
+2. 39 个错误全部是客户端 `ReadTimeout`，不是服务端 OOM/Traceback；服务端日志未发现对应崩溃关键字。
+3. 错误分布：24 个 turn1，15 个 turn2+；按 decode 节点分布为 decode-0 15、decode-1 9、decode-3 7、decode-4 5、decode-5 3。
+4. 8 次 `/open_session` 400 已被 close+retry 兜住，并写入 `structural/session-lifecycle.jsonl`，没有形成 HTTP 400 error row。
+5. 长尾 drain 明显：PD 约 60min 完成，KVC 约 40min 完成；二者都远超 889s trace duration。KVC 在 900s 时已完成 490/626，而 PD 只完成 283/626，说明 KVC 中段吞吐更好，但最后几十个 large-append fallback 仍然拖尾。
+6. direct-to-D 覆盖率为 286/626 = 45.7%，低于样本静态 direct-eligible turn2+ ratio 91.1%。缺口主要来自 D session/residency capacity、large append session cap、reseed/fallback。
+
+当前判断：
+
+- 如果只看 successful paired request，multi-turn-only workload 上 KVC 相对 PD-disaggregation 已经有很强 E2E 提升，且提升主要来自 direct-to-D session reuse。
+- 如果按系统可靠性看，当前实现还不合格，因为 6.23% timeout 会抵消“稳定系统”的结论。
+- 真实 workload 与 microbenchmark 差距的主要原因不是 KVC fast path 无效，而是 fast path 覆盖率不足、D 侧 resident KV/session admission 压力、large append fallback、以及 seeded/reseed path 的 timeout 稳定性。
--- a/docs/REFACTOR_PLAN_ZH.md
+++ b/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 → 写报告。
--- a/docs/STRUCTURAL_VALIDATION_REPORT_ZH.md
+++ b/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 才是）。
--- a/docs/SWEBENCH_EXPERIMENT_PROGRESS.md
+++ b/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，全量回放非常耗时
--- a/docs/SWEBENCH_EXPERIMENT_RESULTS.md
+++ b/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
--- a/docs/V5_PROFILE_INVESTIGATION_ZH.md
+++ b/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` 拿机器可读输出)。
--- a/outputs/qwen3-30b-tp1-v3-kvaware/exp1_1p7d_kvc_kvaware_summary.json
+++ b/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
+  }
+}
--- a/outputs/qwen3-30b-tp1-v3-kvaware/exp2_2p6d_kvc_kvaware_summary.json
+++ b/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
+  }
+}
--- a/outputs/qwen3-30b-tp1-v3-kvaware/sweep_results.txt
+++ b/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 ===
--- a/outputs/qwen3-30b-tp1-v4-cap16/exp1_1p7d_kvc_cap16_summary.json
+++ b/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
+  }
+}
--- a/outputs/qwen3-30b-tp1-v4-cap16/exp2_2p6d_kvc_cap16_summary.json
+++ b/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
+  }
+}
--- a/outputs/qwen3-30b-tp1-v4-cap16/sweep_results.txt
+++ b/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 ===
--- a/pyproject.toml
+++ b/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"
--- a/scripts/analysis/analyze_backpressure_smoke.py
+++ b/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()
--- a/scripts/analysis/analyze_errors.py
+++ b/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}%")
--- a/scripts/analysis/analyze_pool_timeseries.py
+++ b/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()
--- a/scripts/analysis/analyze_v3.py
+++ b/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}%)")
--- a/scripts/analysis/analyze_v4.py
+++ b/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())}")
--- a/scripts/analysis/compare_no_error.py
+++ b/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')}")
--- a/scripts/convert_audit_to_trace.py
+++ b/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]))
--- a/scripts/prepare_real_ali_samples.py
+++ b/scripts/prepare_real_ali_samples.py
@@ -0,0 +1,450 @@
+#!/usr/bin/env python3
+"""Prepare balanced real-Ali trace samples for KVC experiments.
+
+The generic sampler is duration-oriented and can be dominated by one long
+session.  This script keeps real request lengths/timestamps but caps turns per
+session so live sweeps can compare policies on a repeatable multi-session
+workload.
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import statistics
+from collections import defaultdict
+from dataclasses import asdict, dataclass
+from pathlib import Path
+
+from agentic_pd_hybrid.trace import TraceRequest, load_trace
+
+
+@dataclass(frozen=True)
+class SampleSummary:
+    input_trace_path: str
+    output_trace_path: str
+    profile: str
+    request_count: int
+    session_count: int
+    multi_turn_session_count: int
+    turn2plus_count: int
+    direct_eligible_turn2plus_count: int
+    direct_eligible_turn2plus_ratio: float
+    missing_parent_count: int
+    max_sessions: int
+    max_turns_per_session: int
+    start_time_s: float
+    end_time_s: float
+    sampled_duration_s: float
+    rebased_timestamps: bool
+    input_tokens: dict[str, float] | None
+    output_tokens: dict[str, float] | None
+    append_tokens: dict[str, float] | None
+    inter_turn_gap_s: dict[str, float] | None
+    overlap_ratio: dict[str, float] | None
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--trace", type=Path, required=True)
+    parser.add_argument("--output-root", type=Path, required=True)
+    parser.add_argument("--max-sessions", type=int, default=64)
+    parser.add_argument("--max-turns-per-session", type=int, default=12)
+    parser.add_argument("--start-time-s", type=float, default=0.0)
+    parser.add_argument(
+        "--window-duration-s",
+        type=float,
+        default=None,
+        help=(
+            "If set, also write continuous-window samples that keep only requests "
+            "inside [start-time, start-time + window-duration]."
+        ),
+    )
+    parser.add_argument(
+        "--window-target-requests",
+        type=int,
+        default=None,
+        help=(
+            "For continuous-window samples, select whole sessions across time "
+            "buckets until at least this many requests are included. This keeps "
+            "the window span while making live runs tractable."
+        ),
+    )
+    parser.add_argument(
+        "--window-buckets",
+        type=int,
+        default=15,
+        help="Number of time buckets used with --window-target-requests.",
+    )
+    parser.add_argument(
+        "--window-min-turns",
+        type=int,
+        default=1,
+        help=(
+            "Minimum number of in-window turns per selected session for "
+            "continuous-window samples."
+        ),
+    )
+    parser.add_argument(
+        "--window-output-name",
+        default="ali-window.jsonl",
+        help="Output filename for the continuous-window sample.",
+    )
+    parser.add_argument(
+        "--max-sampled-duration-s",
+        type=float,
+        default=None,
+        help=(
+            "For balanced profile samples, drop requests after the first selected "
+            "timestamp plus this duration. Use only for quick smoke runs; headline "
+            "runs should preserve the full sampled span."
+        ),
+    )
+    parser.add_argument(
+        "--profiles",
+        nargs="+",
+        default=["representative-mt", "kvc-fit-smallappend"],
+        choices=["representative-mt", "kvc-fit-smallappend"],
+    )
+    parser.add_argument(
+        "--no-rebase-timestamps",
+        action="store_true",
+        help="Keep original timestamps instead of shifting the sample to start at 0.",
+    )
+    args = parser.parse_args()
+
+    requests = load_trace(args.trace)
+    sessions: dict[str, list[TraceRequest]] = defaultdict(list)
+    for request in requests:
+        sessions[request.session_id].append(request)
+
+    args.output_root.mkdir(parents=True, exist_ok=True)
+    if args.window_duration_s is not None:
+        if args.window_target_requests is None:
+            selected = _select_window(
+                requests=requests,
+                start_time_s=args.start_time_s,
+                window_duration_s=args.window_duration_s,
+            )
+            profile = "window"
+        else:
+            selected = _select_window_session_sample(
+                sessions=sessions,
+                start_time_s=args.start_time_s,
+                window_duration_s=args.window_duration_s,
+                target_requests=args.window_target_requests,
+                bucket_count=args.window_buckets,
+                min_turns=args.window_min_turns,
+            )
+            profile = (
+                "window-session-sample"
+                if args.window_min_turns <= 1
+                else f"window-session-sample-min{args.window_min_turns}turns"
+            )
+        output_path = args.output_root / args.window_output_name
+        summary = _write_sample(
+            selected=selected,
+            input_trace_path=args.trace,
+            output_path=output_path,
+            profile=profile,
+            max_sessions=args.max_sessions,
+            max_turns_per_session=args.max_turns_per_session,
+            rebase_timestamps=not args.no_rebase_timestamps,
+        )
+        print(
+            f"window: wrote {summary.request_count} requests from "
+            f"{summary.session_count} sessions to {output_path}"
+        )
+
+    for profile in args.profiles:
+        selected = _select_profile(
+            sessions=sessions,
+            profile=profile,
+            start_time_s=args.start_time_s,
+            max_sessions=args.max_sessions,
+            max_turns_per_session=args.max_turns_per_session,
+            max_sampled_duration_s=args.max_sampled_duration_s,
+        )
+        output_path = args.output_root / f"ali-{profile}.jsonl"
+        summary = _write_sample(
+            selected=selected,
+            input_trace_path=args.trace,
+            output_path=output_path,
+            profile=profile,
+            max_sessions=args.max_sessions,
+            max_turns_per_session=args.max_turns_per_session,
+            rebase_timestamps=not args.no_rebase_timestamps,
+        )
+        print(
+            f"{profile}: wrote {summary.request_count} requests from "
+            f"{summary.session_count} sessions to {output_path}"
+        )
+
+
+def _select_profile(
+    *,
+    sessions: dict[str, list[TraceRequest]],
+    profile: str,
+    start_time_s: float,
+    max_sessions: int,
+    max_turns_per_session: int,
+    max_sampled_duration_s: float | None,
+) -> list[TraceRequest]:
+    eligible: list[list[TraceRequest]] = []
+    for session_requests in sessions.values():
+        ordered = _ordered(session_requests)
+        if len(ordered) < 2:
+            continue
+        if ordered[0].timestamp_s < start_time_s:
+            continue
+        if profile == "kvc-fit-smallappend" and not _is_kvc_fit_smallappend(ordered):
+            continue
+        eligible.append(ordered[:max_turns_per_session])
+
+    eligible.sort(key=lambda items: (items[0].timestamp_s, items[0].session_id))
+    selected_sessions = eligible[:max_sessions]
+    selected = [request for items in selected_sessions for request in items]
+    selected.sort(key=lambda request: (request.timestamp_s, request.chat_id))
+    if selected and max_sampled_duration_s is not None:
+        first_ts = selected[0].timestamp_s
+        end_ts = first_ts + max_sampled_duration_s
+        selected = [
+            request for request in selected if request.timestamp_s <= end_ts
+        ]
+    return selected
+
+
+def _select_window(
+    *,
+    requests: list[TraceRequest],
+    start_time_s: float,
+    window_duration_s: float,
+) -> list[TraceRequest]:
+    end_time_s = start_time_s + window_duration_s
+    selected = [
+        request
+        for request in requests
+        if start_time_s <= request.timestamp_s <= end_time_s
+    ]
+    selected.sort(key=lambda request: (request.timestamp_s, request.chat_id))
+    return selected
+
+
+def _select_window_session_sample(
+    *,
+    sessions: dict[str, list[TraceRequest]],
+    start_time_s: float,
+    window_duration_s: float,
+    target_requests: int,
+    bucket_count: int,
+    min_turns: int,
+) -> list[TraceRequest]:
+    if target_requests <= 0:
+        raise ValueError("--window-target-requests must be positive")
+    if bucket_count <= 0:
+        raise ValueError("--window-buckets must be positive")
+    if min_turns <= 0:
+        raise ValueError("--window-min-turns must be positive")
+
+    end_time_s = start_time_s + window_duration_s
+    bucket_width_s = window_duration_s / bucket_count
+    buckets: list[list[list[TraceRequest]]] = [[] for _ in range(bucket_count)]
+    for session_requests in sessions.values():
+        ordered = _ordered(session_requests)
+        if not ordered:
+            continue
+        first = ordered[0]
+        if first.timestamp_s < start_time_s or first.timestamp_s > end_time_s:
+            continue
+        in_window = [
+            request
+            for request in ordered
+            if start_time_s <= request.timestamp_s <= end_time_s
+        ]
+        if len(in_window) < min_turns:
+            continue
+        bucket_index = min(
+            bucket_count - 1,
+            int((first.timestamp_s - start_time_s) / bucket_width_s),
+        )
+        buckets[bucket_index].append(in_window)
+
+    for bucket in buckets:
+        bucket.sort(key=lambda items: (items[0].timestamp_s, items[0].session_id))
+
+    selected_sessions: list[list[TraceRequest]] = []
+    selected_count = 0
+    positions = [0 for _ in range(bucket_count)]
+    while selected_count < target_requests:
+        progressed = False
+        for index, bucket in enumerate(buckets):
+            if positions[index] >= len(bucket):
+                continue
+            session_requests = bucket[positions[index]]
+            positions[index] += 1
+            selected_sessions.append(session_requests)
+            selected_count += len(session_requests)
+            progressed = True
+            if selected_count >= target_requests:
+                break
+        if not progressed:
+            break
+
+    selected = [request for items in selected_sessions for request in items]
+    selected.sort(key=lambda request: (request.timestamp_s, request.chat_id))
+    if len(selected) < target_requests:
+        raise ValueError(
+            f"window session sample selected only {len(selected)} requests; "
+            f"target was {target_requests}"
+        )
+    return selected
+
+
+def _is_kvc_fit_smallappend(session_requests: list[TraceRequest]) -> bool:
+    initial = session_requests[0]
+    if initial.input_length < 2048 or initial.input_length > 16000:
+        return False
+    for request in session_requests:
+        if request.output_length > 2048:
+            return False
+    for previous, current in zip(session_requests, session_requests[1:], strict=False):
+        append_tokens = current.input_length - (
+            previous.input_length + previous.output_length
+        )
+        if append_tokens <= 0 or append_tokens > 2048:
+            return False
+        if _overlap_ratio(previous, current) < 0.75:
+            return False
+    return True
+
+
+def _write_sample(
+    *,
+    selected: list[TraceRequest],
+    input_trace_path: Path,
+    output_path: Path,
+    profile: str,
+    max_sessions: int,
+    max_turns_per_session: int,
+    rebase_timestamps: bool,
+) -> SampleSummary:
+    if not selected:
+        raise ValueError(f"profile {profile!r} selected no requests")
+
+    first_ts = selected[0].timestamp_s
+    output_path.parent.mkdir(parents=True, exist_ok=True)
+    with output_path.open("w", encoding="utf-8") as handle:
+        for request in selected:
+            timestamp = request.timestamp_s - first_ts if rebase_timestamps else request.timestamp_s
+            payload = {
+                "chat_id": request.chat_id,
+                "parent_chat_id": request.parent_chat_id,
+                "timestamp": round(timestamp, 6),
+                "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")
+
+    sessions = defaultdict(list)
+    for request in selected:
+        sessions[request.session_id].append(request)
+
+    selected_chat_ids = {request.chat_id for request in selected}
+    missing_parent_count = sum(
+        1
+        for request in selected
+        if request.parent_chat_id >= 0 and request.parent_chat_id not in selected_chat_ids
+    )
+    append_values: list[float] = []
+    gap_values: list[float] = []
+    overlap_values: list[float] = []
+    direct_eligible_count = 0
+    for session_requests in sessions.values():
+        ordered = _ordered(session_requests)
+        for previous, current in zip(ordered, ordered[1:], strict=False):
+            append_tokens = current.input_length - (
+                previous.input_length + previous.output_length
+            )
+            overlap_ratio = _overlap_ratio(previous, current)
+            append_values.append(float(append_tokens))
+            gap_values.append(float(current.timestamp_s - previous.timestamp_s))
+            overlap_values.append(overlap_ratio)
+            if append_tokens > 0 and append_tokens <= 2048 and overlap_ratio > 0:
+                direct_eligible_count += 1
+
+    turn2plus_count = sum(max(0, len(items) - 1) for items in sessions.values())
+
+    start = min(request.timestamp_s for request in selected)
+    end = max(request.timestamp_s for request in selected)
+    summary = SampleSummary(
+        input_trace_path=str(input_trace_path),
+        output_trace_path=str(output_path),
+        profile=profile,
+        request_count=len(selected),
+        session_count=len(sessions),
+        multi_turn_session_count=sum(1 for items in sessions.values() if len(items) > 1),
+        turn2plus_count=turn2plus_count,
+        direct_eligible_turn2plus_count=direct_eligible_count,
+        direct_eligible_turn2plus_ratio=(
+            direct_eligible_count / turn2plus_count if turn2plus_count else 0.0
+        ),
+        missing_parent_count=missing_parent_count,
+        max_sessions=max_sessions,
+        max_turns_per_session=max_turns_per_session,
+        start_time_s=0.0 if rebase_timestamps else start,
+        end_time_s=end - start if rebase_timestamps else end,
+        sampled_duration_s=end - start,
+        rebased_timestamps=rebase_timestamps,
+        input_tokens=_stats([float(request.input_length) for request in selected]),
+        output_tokens=_stats([float(request.output_length) for request in selected]),
+        append_tokens=_stats(append_values),
+        inter_turn_gap_s=_stats(gap_values),
+        overlap_ratio=_stats(overlap_values),
+    )
+    with output_path.with_suffix(output_path.suffix + ".summary.json").open(
+        "w", encoding="utf-8"
+    ) as handle:
+        json.dump(asdict(summary), handle, indent=2, sort_keys=True)
+    return summary
+
+
+def _ordered(session_requests: list[TraceRequest]) -> list[TraceRequest]:
+    return sorted(
+        session_requests,
+        key=lambda request: (request.timestamp_s, request.turn_id, request.chat_id),
+    )
+
+
+def _overlap_ratio(previous: TraceRequest, current: TraceRequest) -> float:
+    if not current.hash_ids:
+        return 0.0
+    previous_blocks = set(previous.hash_ids)
+    overlap = sum(1 for block in current.hash_ids if block in previous_blocks)
+    return overlap / len(current.hash_ids)
+
+
+def _stats(values: list[float]) -> dict[str, float] | None:
+    if not values:
+        return None
+    ordered = sorted(values)
+    return {
+        "count": float(len(ordered)),
+        "mean": statistics.fmean(ordered),
+        "min": ordered[0],
+        "p50": _percentile(ordered, 0.50),
+        "p90": _percentile(ordered, 0.90),
+        "p99": _percentile(ordered, 0.99),
+        "max": ordered[-1],
+    }
+
+
+def _percentile(sorted_values: list[float], percentile: float) -> float:
+    if len(sorted_values) == 1:
+        return sorted_values[0]
+    return sorted_values[round((len(sorted_values) - 1) * percentile)]
+
+
+if __name__ == "__main__":
+    main()
--- a/scripts/run_all_experiments.sh
+++ b/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 ==="
--- a/scripts/run_exp_a_pd_disagg.sh
+++ b/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
--- a/scripts/run_exp_b1_dp_colo_rr.sh
+++ b/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
--- a/scripts/run_exp_b2_dp_colo_cache_aware.sh
+++ b/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
--- a/scripts/run_exp_b_pd_colo.sh
+++ b/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
--- a/scripts/run_exp_c_kvcache_centric.sh
+++ b/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
--- a/scripts/smoke_test.sh
+++ b/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
--- a/scripts/sweep_backpressure_smoke.sh
+++ b/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"
--- a/scripts/sweep_kvc_qwen3_30b.sh
+++ b/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)"
--- a/scripts/sweep_real_ali_kvc.sh
+++ b/scripts/sweep_real_ali_kvc.sh
@@ -0,0 +1,170 @@
+#!/usr/bin/env bash
+# Real Ali workload sweep for KVC pd-hybrid.
+#
+# This script expects a prebuilt sample trace and replays it exactly for every
+# mechanism.  It intentionally keeps pool polling disabled for performance runs.
+set -euo pipefail
+
+REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+cd "$REPO_ROOT"
+
+MODEL=${MODEL:-/home/admin/cpfs/wjh/models/Qwen/Qwen3-Coder-30B-A3B-Instruct}
+TRACE=${TRACE:-outputs/real-ali-kvc-iter/samples-balanced/ali-kvc-fit-smallappend.jsonl}
+OUT_ROOT=${OUT_ROOT:-outputs/real-ali-kvc-iter/runs}
+TIME_SCALE=${TIME_SCALE:-1}
+CONCURRENCY=${CONCURRENCY:-32}
+REQUEST_TIMEOUT_S=${REQUEST_TIMEOUT_S:-300}
+STACK_TIMEOUT_S=${STACK_TIMEOUT_S:-1200}
+RUNS=${RUNS:-"dp kvc_bp"}
+EXTRA_SERVER_ARGS=${EXTRA_SERVER_ARGS:-}
+PREFILL_EXTRA_SERVER_ARGS=${PREFILL_EXTRA_SERVER_ARGS:-}
+DECODE_EXTRA_SERVER_ARGS=${DECODE_EXTRA_SERVER_ARGS:-}
+KVC_SEED_MIN_TURN_ID=${KVC_SEED_MIN_TURN_ID:-1}
+KVC_SEED_ONLY_MULTITURN=${KVC_SEED_ONLY_MULTITURN:-0}
+
+mkdir -p "$OUT_ROOT"
+LOG="$OUT_ROOT/sweep.log"
+
+log() {
+  echo "[$(date '+%F %T')] $*" | tee -a "$LOG"
+}
+
+common_args=(
+  --trace "$TRACE"
+  --model-path "$MODEL"
+  --output-root "$OUT_ROOT"
+  --use-trace-as-sample
+  --time-scale "$TIME_SCALE"
+  --concurrency-limit "$CONCURRENCY"
+  --timeout-s "$STACK_TIMEOUT_S"
+  --request-timeout-s "$REQUEST_TIMEOUT_S"
+)
+if [[ -n "$EXTRA_SERVER_ARGS" ]]; then
+  common_args+=(--extra-server-args "$EXTRA_SERVER_ARGS")
+fi
+if [[ -n "$PREFILL_EXTRA_SERVER_ARGS" ]]; then
+  common_args+=(--prefill-extra-server-args "$PREFILL_EXTRA_SERVER_ARGS")
+fi
+if [[ -n "$DECODE_EXTRA_SERVER_ARGS" ]]; then
+  common_args+=(--decode-extra-server-args "$DECODE_EXTRA_SERVER_ARGS")
+fi
+
+kvc_args=(
+  "${common_args[@]}"
+  --mechanism kvcache-centric
+  --policy kv-aware
+  --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
+  --kvcache-admission-mode worker
+  --kvcache-seed-min-turn-id "$KVC_SEED_MIN_TURN_ID"
+  --kvcache-seed-max-inflight-decode -1
+  --kvcache-prefill-backup-policy release-after-transfer
+  --kvcache-prefill-priority-eviction
+)
+if [[ "$KVC_SEED_ONLY_MULTITURN" == "1" ]]; then
+  kvc_args+=(--kvcache-seed-only-multiturn-sessions)
+fi
+
+run_dp() {
+  log "=== DP cache-aware baseline: 8 direct workers ==="
+  uv run agentic-pd-hybrid benchmark-live \
+    "${common_args[@]}" \
+    --mechanism pd-colo \
+    --policy kv-aware \
+    --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
+}
+
+run_pd_disagg() {
+  log "=== PD-disaggregation baseline: 2P6D ==="
+  uv run agentic-pd-hybrid benchmark-live \
+    "${common_args[@]}" \
+    --mechanism pd-disaggregation \
+    --policy kv-aware \
+    --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
+}
+
+run_pd_sticky() {
+  log "=== PD-disaggregation sticky baseline: 2P6D ==="
+  uv run agentic-pd-hybrid benchmark-live \
+    "${common_args[@]}" \
+    --mechanism pd-disaggregation \
+    --policy sticky \
+    --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
+}
+
+run_kvc() {
+  log "=== KVC baseline: 2P6D worker admission, no backpressure ==="
+  uv run agentic-pd-hybrid benchmark-live "${kvc_args[@]}"
+}
+
+run_kvc_bp() {
+  log "=== KVC candidate: 2P6D worker admission + backpressure ==="
+  uv run agentic-pd-hybrid benchmark-live \
+    "${kvc_args[@]}" \
+    --enable-backpressure \
+    --backpressure-max-pause-s 2.0
+}
+
+summarize_latest() {
+  log "=== Latest summaries ==="
+  find "$OUT_ROOT" -maxdepth 2 -name 'request-metrics.jsonl.summary.json' -print \
+    | sort \
+    | while read -r summary; do
+        python - "$summary" <<'PY'
+import json, sys
+p=sys.argv[1]
+d=json.load(open(p))
+lat=d.get("latency_stats_s") or {}
+tt=d.get("ttft_stats_s") or {}
+em=d.get("execution_modes") or {}
+print(p)
+print("  reqs", d.get("request_count"), "errors", d.get("error_count"), "trunc", d.get("truncated_request_count"))
+print("  lat mean/p50/p90/p99", lat.get("mean"), lat.get("p50"), lat.get("p90"), lat.get("p99"))
+print("  ttft mean/p50/p90", tt.get("mean"), tt.get("p50"), tt.get("p90"))
+print("  modes", em)
+PY
+      done | tee -a "$LOG"
+}
+
+log "Trace: $TRACE"
+log "Model: $MODEL"
+log "Runs: $RUNS | time-scale=$TIME_SCALE concurrency=$CONCURRENCY | kvc-seed-min-turn-id=$KVC_SEED_MIN_TURN_ID | kvc-seed-only-multiturn=$KVC_SEED_ONLY_MULTITURN"
+
+for run in $RUNS; do
+  case "$run" in
+    dp) run_dp ;;
+    pd) run_pd_disagg ;;
+    pd_sticky) run_pd_sticky ;;
+    kvc) run_kvc ;;
+    kvc_bp) run_kvc_bp ;;
+    *) log "Unknown run name: $run"; exit 2 ;;
+  esac
+done
+
+summarize_latest
+log "DONE"
--- a/scripts/sweep_tp1_configs.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v2_fixed.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v3_kvaware.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v4_cap16.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v5_baseline_rerun_exp2.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v5_optD.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v5_optD_profile.sh
+++ b/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 ==="
--- a/scripts/sweep_tp1_v6_p1_profile.sh
+++ b/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 ==="
--- a/src/agentic_pd_hybrid/init.py
+++ b/src/agentic_pd_hybrid/init.py
@@ -0,0 +1,12 @@
+"""Agentic PD hybrid prototype."""
+
+__all__ = [
+    "cli",
+    "launcher",
+    "metrics",
+    "microbench",
+    "policies",
+    "replay",
+    "topology",
+    "trace",
+]
--- a/src/agentic_pd_hybrid/benchmark.py
+++ b/src/agentic_pd_hybrid/benchmark.py
@@ -0,0 +1,380 @@
+from __future__ import annotations
+
+import asyncio
+import json
+import signal
+import shutil
+from collections import Counter
+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,
+    SessionSampleSummary,
+    sample_trace_sessions,
+)
+from agentic_pd_hybrid.stack import ManagedPdStack, launch_pd_stack
+from agentic_pd_hybrid.trace import load_trace
+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
+    progress_interval_s: float = 30.0
+    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
+    use_trace_as_sample: bool = False
+    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"
+    if config.use_trace_as_sample:
+        shutil.copyfile(config.trace_path, sampled_trace_path)
+        sample_summary = _summarize_trace_sample(
+            input_trace_path=config.trace_path,
+            sampled_trace_path=sampled_trace_path,
+            profile=config.sample_profile,
+            session_sample_rate=config.session_sample_rate,
+            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,
+        )
+    else:
+        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,
+            progress_interval_s=config.progress_interval_s,
+        )
+        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,
+                "progress_interval_s": config.progress_interval_s,
+                "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,
+                "use_trace_as_sample": config.use_trace_as_sample,
+                "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"
+
+
+def _summarize_trace_sample(
+    *,
+    input_trace_path: Path,
+    sampled_trace_path: Path,
+    profile: str,
+    session_sample_rate: float,
+    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,
+) -> SessionSampleSummary:
+    requests = load_trace(sampled_trace_path)
+    if not requests:
+        raise ValueError(f"Trace sample is empty: {sampled_trace_path}")
+    session_turns = Counter(request.session_id for request in requests)
+    start_time_s = requests[0].timestamp_s
+    end_time_s = requests[-1].timestamp_s
+    return SessionSampleSummary(
+        input_trace_path=str(input_trace_path),
+        output_trace_path=str(sampled_trace_path),
+        request_count=len(requests),
+        session_count=len(session_turns),
+        multi_turn_session_count=sum(1 for turns in session_turns.values() if turns > 1),
+        start_time_s=start_time_s,
+        end_time_s=end_time_s,
+        sampled_duration_s=end_time_s - start_time_s,
+        session_sample_rate=session_sample_rate,
+        min_turns=min_turns,
+        profile=profile,
+        min_initial_input_tokens=min_initial_input_tokens,
+        max_initial_input_tokens=max_initial_input_tokens,
+        max_append_input_tokens=max_append_input_tokens,
+        max_output_tokens=max_output_tokens,
+        min_overlap_ratio=min_overlap_ratio,
+        mean_append_input_tokens=None,
+        mean_turn_overlap_ratio=None,
+    )
--- a/src/agentic_pd_hybrid/cli.py
+++ b/src/agentic_pd_hybrid/cli.py
@@ -0,0 +1,890 @@
+from __future__ import annotations
+
+import argparse
+import asyncio
+import shlex
+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(
+        "--progress-interval-s",
+        type=float,
+        default=30.0,
+        help=(
+            "Write client-side replay progress to <output_dir>/replay-progress.jsonl "
+            "every N seconds. 0 disables the heartbeat."
+        ),
+    )
+
+    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(
+        "--progress-interval-s",
+        type=float,
+        default=30.0,
+        help=(
+            "Write client-side replay progress to <run_dir>/replay-progress.jsonl "
+            "every N seconds. 0 disables the heartbeat."
+        ),
+    )
+    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)
+    benchmark.add_argument(
+        "--use-trace-as-sample",
+        action="store_true",
+        help=(
+            "Replay the provided --trace exactly instead of sampling sessions into "
+            "a new trace. Use this for prebuilt real-workload samples."
+        ),
+    )
+
+    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,
+            progress_interval_s=args.progress_interval_s,
+        )
+        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,
+                progress_interval_s=args.progress_interval_s,
+                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,
+                use_trace_as_sample=args.use_trace_as_sample,
+                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",
+    )
+    parser.add_argument(
+        "--extra-server-args",
+        default="",
+        help="Extra arguments appended to every sglang.launch_server command.",
+    )
+    parser.add_argument(
+        "--prefill-extra-server-args",
+        default="",
+        help="Extra arguments appended only to prefill launch_server commands.",
+    )
+    parser.add_argument(
+        "--decode-extra-server-args",
+        default="",
+        help="Extra arguments appended only to decode launch_server commands.",
+    )
+    parser.add_argument(
+        "--direct-extra-server-args",
+        default="",
+        help="Extra arguments appended only to direct launch_server commands.",
+    )
+
+
+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,
+        extra_server_args=tuple(shlex.split(args.extra_server_args)),
+        prefill_extra_server_args=tuple(shlex.split(args.prefill_extra_server_args)),
+        decode_extra_server_args=tuple(shlex.split(args.decode_extra_server_args)),
+        direct_extra_server_args=(
+            "--enable-streaming-session",
+            *tuple(shlex.split(args.direct_extra_server_args)),
+        ),
+    )
+
+
+if __name__ == "__main__":
+    main()
--- a/src/agentic_pd_hybrid/launcher.py
+++ b/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
--- a/src/agentic_pd_hybrid/metrics.py
+++ b/src/agentic_pd_hybrid/metrics.py
@@ -0,0 +1,207 @@
+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 write_summary_json(
+    path: Path,
+    rows: list[RequestMetrics],
+    *,
+    trace_path: Path,
+    router_url: str | None,
+) -> None:
+    latencies = [row.latency_s for row in rows if row.latency_s is not None]
+    ttfts = [row.ttft_s for row in rows if row.ttft_s is not None]
+    tpots = [row.tpot_s for row in rows 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),
+        "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]
--- a/src/agentic_pd_hybrid/microbench.py
+++ b/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
--- a/src/agentic_pd_hybrid/pd_router.py
+++ b/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()
--- a/src/agentic_pd_hybrid/policies.py
+++ b/src/agentic_pd_hybrid/policies.py
@@ -0,0 +1,235 @@
+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)
+
+    @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 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
+
+    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] | None = None
+        for worker in topology.route_workers:
+            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
+
+        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) -> 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()
+    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)
--- a/src/agentic_pd_hybrid/profile.py
+++ b/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"
--- a/src/agentic_pd_hybrid/replay.py
+++ b/src/agentic_pd_hybrid/replay.py
--- a/src/agentic_pd_hybrid/sampling.py
+++ b/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
--- a/src/agentic_pd_hybrid/stack.py
+++ b/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})")
--- a/src/agentic_pd_hybrid/topology.py
+++ b/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,
+    )
--- a/src/agentic_pd_hybrid/trace.py
+++ b/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
--- a/src/agentic_pd_hybrid/trace_profiles.py
+++ b/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)
--- a/third_party/sglang/.claude/skills/add-jit-kernel/SKILL.md
+++ b/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
+```
--- a/third_party/sglang/.claude/skills/add-sgl-kernel/SKILL.md
+++ b/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
+```
--- a/third_party/sglang/.claude/skills/ci-workflow-guide/SKILL.md
+++ b/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`.
--- a/third_party/sglang/.claude/skills/debug-cuda-crash/SKILL.md
+++ b/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.
--- a/third_party/sglang/.claude/skills/generate-profile/SKILL.md
+++ b/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.
--- a/third_party/sglang/.claude/skills/sglang-bisect-ci-regression/SKILL.md
+++ b/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.
--- a/third_party/sglang/.claude/skills/write-sglang-test/SKILL.md
+++ b/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)
--- a/third_party/sglang/.codespellrc
+++ b/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
--- a/third_party/sglang/.coveragerc
+++ b/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
--- a/third_party/sglang/.devcontainer/Dockerfile
+++ b/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
--- a/third_party/sglang/.devcontainer/devcontainer.json
+++ b/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"
+}
--- a/third_party/sglang/.dockerignore
+++ b/third_party/sglang/.dockerignore
@@ -0,0 +1 @@
+.gitignore
--- a/third_party/sglang/.github/CI_PERMISSIONS.json
+++ b/third_party/sglang/.github/CI_PERMISSIONS.json
--- a/third_party/sglang/.github/CODEOWNERS
+++ b/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
--- a/third_party/sglang/.github/FOLDER_README.md
+++ b/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.
--- a/third_party/sglang/.github/ISSUE_TEMPLATE/1-bug-report.yml
+++ b/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
--- a/third_party/sglang/.github/ISSUE_TEMPLATE/2-feature-request.yml
+++ b/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.
--- a/third_party/sglang/.github/MAINTAINER.md
+++ b/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.
--- a/third_party/sglang/.github/actions/check-maintenance/action.yml
+++ b/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
--- a/third_party/sglang/.github/actions/check-stage-health/action.yml
+++ b/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(', ')}`);
+          }
--- a/third_party/sglang/.github/actions/upload-cuda-coredumps/action.yml
+++ b/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' }}"
--- a/third_party/sglang/.github/actions/wait-for-jobs/action.yml
+++ b/third_party/sglang/.github/actions/wait-for-jobs/action.yml
@@ -0,0 +1,177 @@
+name: Wait for Jobs
+description: Poll and wait for specified jobs in the current workflow run to complete
+
+inputs:
+  stage-name:
+    description: 'Human-readable stage name for log messages (e.g. "stage-a")'
+    required: true
+  jobs:
+    description: |
+      JSON array of job specs to wait for. Each element is either:
+        - a string: exact job name (e.g. "stage-a-test-1-gpu-small")
+        - an object { "prefix": "...", "expected_count": N }: for matrix jobs
+    required: true
+  max-wait-minutes:
+    description: 'Maximum time to wait before timing out'
+    required: false
+    default: '240'
+  poll-interval-seconds:
+    description: 'Seconds between polling attempts'
+    required: false
+    default: '60'
+  github-token:
+    description: 'GitHub token for API calls'
+    required: false
+    default: ${{ github.token }}
+
+outputs:
+  result:
+    description: 'Overall result: success, failure, or timeout'
+    value: ${{ steps.wait.outputs.result }}
+
+runs:
+  using: composite
+  steps:
+    - name: Wait for jobs to complete
+      id: wait
+      uses: actions/github-script@v7
+      env:
+        INPUT_STAGE_NAME: ${{ inputs.stage-name }}
+        INPUT_JOBS: ${{ inputs.jobs }}
+        INPUT_MAX_WAIT_MINUTES: ${{ inputs.max-wait-minutes }}
+        INPUT_POLL_INTERVAL_SECONDS: ${{ inputs.poll-interval-seconds }}
+      with:
+        github-token: ${{ inputs.github-token }}
+        script: |
+          const stageName = process.env.INPUT_STAGE_NAME;
+          const jobSpecs = JSON.parse(process.env.INPUT_JOBS);
+          const maxWaitMinutes = parseInt(process.env.INPUT_MAX_WAIT_MINUTES);
+          const pollIntervalSeconds = parseInt(process.env.INPUT_POLL_INTERVAL_SECONDS);
+          const maxAttempts = (maxWaitMinutes * 60) / pollIntervalSeconds;
+
+          // Normalize job specs into a uniform format
+          const normalizedSpecs = jobSpecs.map(spec => {
+            if (typeof spec === 'string') {
+              return { prefix: spec, expected_count: 1, exact: true };
+            }
+            return { ...spec, exact: false };
+          });
+
+          const totalExpectedJobs = normalizedSpecs.reduce((sum, s) => sum + s.expected_count, 0);
+
+          const matchesSpec = (jobName, spec) => {
+            if (spec.exact) {
+              return jobName === spec.prefix;
+            }
+            return jobName === spec.prefix || jobName.startsWith(spec.prefix + ' (');
+          };
+
+          // Use ETag conditional requests to avoid consuming rate limit when nothing changed.
+          // GitHub returns 304 Not Modified for unchanged data, which is FREE (no rate limit cost).
+          let lastEtag = '';
+          let lastJobs = null;
+          let apiCalls = 0;
+          let cachedCalls = 0;
+
+          async function fetchJobs() {
+            const url = `GET /repos/{owner}/{repo}/actions/runs/{run_id}/jobs`;
+            const params = {
+              owner: context.repo.owner,
+              repo: context.repo.repo,
+              run_id: context.runId,
+              per_page: 100,
+              headers: {},
+            };
+            if (lastEtag) {
+              params.headers['if-none-match'] = lastEtag;
+            }
+
+            try {
+              const response = await github.request(url, params);
+              apiCalls++;
+              const rateRemaining = response.headers['x-ratelimit-remaining'] || '?';
+              const rateLimit = response.headers['x-ratelimit-limit'] || '?';
+              console.log(`[rate-limit] ${rateRemaining}/${rateLimit} remaining (ETag: ${lastEtag ? 'sent' : 'none'}) | this session: ${apiCalls} paid, ${cachedCalls} free`);
+              lastEtag = response.headers.etag || '';
+              const jobs = response.data.jobs;
+
+              // Handle pagination if >100 jobs
+              // ETag only covers page 1, so invalidate it to avoid stale cache
+              // when later pages change but page 1 doesn't.
+              if (response.data.total_count > 100) {
+                lastEtag = '';
+                for (let page = 2; page <= Math.ceil(response.data.total_count / 100); page++) {
+                  const { data: pageData } = await github.request(url, {
+                    ...params,
+                    page,
+                    headers: {},
+                  });
+                  jobs.push(...pageData.jobs);
+                }
+              }
+
+              lastJobs = jobs;
+              return { jobs, cached: false };
+            } catch (err) {
+              if (err.status === 304 && lastJobs) {
+                cachedCalls++;
+                console.log(`[rate-limit] 304 Not Modified | this session: ${apiCalls} paid, ${cachedCalls} free`);
+                return { jobs: lastJobs, cached: true };
+              }
+              throw err;
+            }
+          }
+
+          for (let attempt = 0; attempt < maxAttempts; attempt++) {
+            const { jobs, cached } = await fetchJobs();
+
+            let allCompleted = true;
+            let failedJobs = [];
+            let completedCount = 0;
+            let totalCount = 0;
+
+            for (const spec of normalizedSpecs) {
+              const matchingJobs = jobs.filter(job => matchesSpec(job.name, spec));
+
+              for (const job of matchingJobs) {
+                totalCount++;
+                if (!cached) {
+                  console.log(`${job.name}: status=${job.status}, conclusion=${job.conclusion}`);
+                }
+
+                if (job.status === 'completed') {
+                  completedCount++;
+                  if (job.conclusion !== 'success' && job.conclusion !== 'skipped') {
+                    failedJobs.push(job.name);
+                  }
+                } else {
+                  allCompleted = false;
+                }
+              }
+
+              if (matchingJobs.length < spec.expected_count) {
+                console.log(`${spec.prefix}: found ${matchingJobs.length}/${spec.expected_count} jobs (waiting for more)`);
+                allCompleted = false;
+              }
+            }
+
+            console.log(`[${stageName}] Progress: ${completedCount}/${totalCount} jobs completed (expected ${totalExpectedJobs})${cached ? ' (cached, no rate limit cost)' : ''}`);
+
+            // Fail fast if any jobs failed
+            if (failedJobs.length > 0) {
+              core.setOutput('result', 'failure');
+              core.setFailed(`${stageName} jobs failed: ${failedJobs.join(', ')}`);
+              return;
+            }
+
+            if (allCompleted && totalCount >= totalExpectedJobs) {
+              core.setOutput('result', 'success');
+              return;
+            }
+
+            console.log(`Waiting ${pollIntervalSeconds}s... (attempt ${attempt + 1}/${maxAttempts})`);
+            await new Promise(resolve => setTimeout(resolve, pollIntervalSeconds * 1000));
+          }
+
+          core.setFailed(`Timeout waiting for ${stageName} jobs`);
+          core.setOutput('result', 'timeout');
--- a/third_party/sglang/.github/audit_permission.py
+++ b/third_party/sglang/.github/audit_permission.py
@@ -0,0 +1,411 @@
+"""
+Audit GitHub repository collaborators with elevated access.
+
+This script will:
+1. Fetch all collaborators with write permission to this repo.
+2. Show their github username, Nickname and the role (e.g., admin, maintain,
+   custom org role, write, triage).
+3. Show their last activity related to this repo (last commit, last issue,
+   last pull request). Put the data in YYYY-MM-DD format. Add a column "last activity date" to the CSV, before the above three breakdown columns.
+4. Show activity on other repos: repos touched via public events in the last 90 days (Push, PR, Issues, etc.). Sort the repos by the number of activities.
+5. Write results to a CSV sorted by the roles (admin, maintain, custom org role, write, triage) and the last activity date (most recent first).
+
+Usage:
+    export GH_TOKEN="your_github_token"
+    python3 audit_permission.py [--output path] [--repo owner/name]
+
+Requires: requests, and a token with permission to list collaborators (push+
+access to the repo).
+"""
+
+from __future__ import annotations
+
+import argparse
+import csv
+import os
+import sys
+import time
+from collections import Counter
+from datetime import datetime, timedelta, timezone
+from typing import Any
+
+try:
+    import requests
+except ImportError:
+    requests = None  # type: ignore
+
+DEFAULT_OWNER = "sgl-project"
+DEFAULT_NAME = "sglang"
+
+HEADERS: dict[str, str] = {}
+
+
+def _request(
+    method: str,
+    url: str,
+    *,
+    params: dict[str, Any] | None = None,
+    max_retries: int = 3,
+) -> requests.Response:
+    if requests is None:
+        raise RuntimeError("Install the requests package: pip install requests")
+    for attempt in range(max_retries):
+        r = requests.request(method, url, headers=HEADERS, params=params, timeout=60)
+        if r.status_code == 403 and "rate limit" in (r.text or "").lower():
+            reset = r.headers.get("X-RateLimit-Reset")
+            wait = 60
+            if reset:
+                try:
+                    wait = max(1, int(reset) - int(time.time()) + 2)
+                except ValueError:
+                    pass
+            print(f"Rate limited; sleeping {wait}s...", file=sys.stderr)
+            time.sleep(min(wait, 3600))
+            continue
+        return r
+    return r
+
+
+def paginate_list(url: str, params: dict[str, Any] | None = None) -> list[Any]:
+    out: list[Any] = []
+    next_url: str | None = url
+    next_params = params
+    while next_url:
+        r = _request("GET", next_url, params=next_params)
+        next_params = None
+        if r.status_code != 200:
+            print(
+                f"Error {r.status_code} GET {next_url}: {r.text[:500]}",
+                file=sys.stderr,
+            )
+            break
+        data = r.json()
+        if isinstance(data, list):
+            out.extend(data)
+        else:
+            break
+        next_url = None
+        link = r.headers.get("Link", "")
+        for part in link.split(", "):
+            if 'rel="next"' in part:
+                start = part.find("<") + 1
+                end = part.find(">")
+                if start > 0 and end > start:
+                    next_url = part[start:end]
+                break
+    return out
+
+
+def collaborator_role(collab: dict[str, Any]) -> str:
+    role_name = collab.get("role_name")
+    if isinstance(role_name, str) and role_name.strip():
+        return role_name.strip()
+    perms = collab.get("permissions") or {}
+    if perms.get("admin"):
+        return "admin"
+    if perms.get("maintain"):
+        return "maintain"
+    if perms.get("push"):
+        return "write"
+    if perms.get("triage"):
+        return "triage"
+    return "read"
+
+
+def has_write_plus(collab: dict[str, Any]) -> bool:
+    perms = collab.get("permissions") or {}
+    return bool(
+        perms.get("admin")
+        or perms.get("maintain")
+        or perms.get("push")
+        or perms.get("triage")
+    )
+
+
+def role_sort_tier(collab: dict[str, Any]) -> int:
+    """Sort order: admin (0), maintain (1), custom org role (2), write (3), triage (4)."""
+    rn = collab.get("role_name")
+    if isinstance(rn, str) and rn.strip():
+        k = rn.strip().lower()
+        if k == "admin":
+            return 0
+        if k == "maintain":
+            return 1
+        if k == "write":
+            return 3
+        if k == "triage":
+            return 4
+        if k == "read":
+            return 5
+        return 2
+    perms = collab.get("permissions") or {}
+    if perms.get("admin"):
+        return 0
+    if perms.get("maintain"):
+        return 1
+    if perms.get("push"):
+        return 3
+    if perms.get("triage"):
+        return 4
+    return 5
+
+
+def fetch_display_name(login: str) -> str:
+    url = f"https://api.github.com/users/{login}"
+    r = _request("GET", url)
+    if r.status_code != 200:
+        return ""
+    data = r.json()
+    if not isinstance(data, dict):
+        return ""
+    n = data.get("name")
+    return n.strip() if isinstance(n, str) else ""
+
+
+def parse_github_ts(s: str) -> datetime | None:
+    if not s:
+        return None
+    s = s.replace("Z", "+00:00")
+    try:
+        return datetime.fromisoformat(s)
+    except ValueError:
+        return None
+
+
+def iso_timestamp_to_ymd(iso: str | None) -> str:
+    if not iso:
+        return ""
+    p = parse_github_ts(iso)
+    if not p:
+        return ""
+    return p.date().isoformat()
+
+
+def max_date_ymd(*iso_dates: str | None) -> str:
+    best: datetime | None = None
+    for d in iso_dates:
+        p = parse_github_ts(d or "")
+        if p and (best is None or p > best):
+            best = p
+    return best.date().isoformat() if best else ""
+
+
+def parse_ymd(s: str) -> datetime | None:
+    if not s:
+        return None
+    try:
+        return datetime.strptime(s, "%Y-%m-%d").replace(tzinfo=timezone.utc)
+    except ValueError:
+        return None
+
+
+def last_commit_date(owner: str, repo: str, login: str) -> str | None:
+    url = f"https://api.github.com/repos/{owner}/{repo}/commits"
+    r = _request("GET", url, params={"author": login, "per_page": 1})
+    if r.status_code != 200:
+        return None
+    data = r.json()
+    if not isinstance(data, list) or not data:
+        return None
+    commit = data[0].get("commit") or {}
+    c = commit.get("committer") or commit.get("author") or {}
+    d = c.get("date")
+    return d if isinstance(d, str) else None
+
+
+def search_repo_item(
+    owner: str, repo: str, login: str, kind: str
+) -> dict[str, Any] | None:
+    q = f"repo:{owner}/{repo} is:{kind} author:{login}"
+    url = "https://api.github.com/search/issues"
+    r = _request(
+        "GET",
+        url,
+        params={"q": q, "sort": "updated", "order": "desc", "per_page": 1},
+    )
+    if r.status_code != 200:
+        return None
+    payload = r.json()
+    items = payload.get("items")
+    if not items:
+        return None
+    return items[0] if isinstance(items[0], dict) else None
+
+
+def last_issue_pr_dates(
+    owner: str, repo: str, login: str
+) -> tuple[str | None, str | None]:
+    issue = search_repo_item(owner, repo, login, "issue")
+    pr = search_repo_item(owner, repo, login, "pr")
+    issue_dt = None
+    pr_dt = None
+    if issue:
+        issue_dt = issue.get("updated_at") or issue.get("created_at")
+        if not isinstance(issue_dt, str):
+            issue_dt = None
+    if pr:
+        pr_dt = pr.get("updated_at") or pr.get("created_at")
+        if not isinstance(pr_dt, str):
+            pr_dt = None
+    return issue_dt, pr_dt
+
+
+def other_repos_activity_column(
+    login: str, owner: str, repo: str, days: int = 90
+) -> str:
+    """Repos other than this one touched in the window, sorted by event count (desc)."""
+    cutoff = datetime.now(timezone.utc) - timedelta(days=days)
+    full = f"{owner}/{repo}"
+    counts: Counter[str] = Counter()
+    url: str | None = f"https://api.github.com/users/{login}/events/public"
+    params: dict[str, Any] = {"per_page": 100}
+
+    while url:
+        r = _request("GET", url, params=params)
+        params = {}
+        if r.status_code != 200:
+            break
+        events = r.json()
+        if not isinstance(events, list):
+            break
+        oldest_in_page: datetime | None = None
+        for ev in events:
+            if not isinstance(ev, dict):
+                continue
+            created = parse_github_ts(ev.get("created_at") or "")
+            if created:
+                if oldest_in_page is None or created < oldest_in_page:
+                    oldest_in_page = created
+            if created and created < cutoff:
+                continue
+            rinfo = ev.get("repo")
+            name = None
+            if isinstance(rinfo, dict):
+                name = rinfo.get("name")
+            if isinstance(name, str) and name and name != full:
+                counts[name] += 1
+        next_url = None
+        link = r.headers.get("Link", "")
+        for part in link.split(", "):
+            if 'rel="next"' in part:
+                s, e = part.find("<") + 1, part.find(">")
+                if s > 0 and e > s:
+                    next_url = part[s:e]
+                break
+        if oldest_in_page and oldest_in_page < cutoff:
+            break
+        url = next_url
+        if not events:
+            break
+
+    ordered = sorted(counts.items(), key=lambda x: (-x[1], x[0]))
+    return ";".join(f"{n}:{c}" for n, c in ordered)
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser(description="Audit repo collaborator permissions.")
+    parser.add_argument(
+        "--repo",
+        default=f"{DEFAULT_OWNER}/{DEFAULT_NAME}",
+        help=f"owner/name (default: {DEFAULT_OWNER}/{DEFAULT_NAME})",
+    )
+    parser.add_argument(
+        "--output",
+        "-o",
+        default=os.path.join(os.path.dirname(__file__), "permission_audit.csv"),
+        help="Output CSV path",
+    )
+    parser.add_argument(
+        "--events-days",
+        type=int,
+        default=90,
+        help="Window for other-repo activity via public events",
+    )
+    args = parser.parse_args()
+
+    if "/" not in args.repo:
+        print("Error: --repo must be owner/name", file=sys.stderr)
+        sys.exit(1)
+    owner, name = args.repo.split("/", 1)
+
+    gh_token = os.getenv("GH_TOKEN")
+    if not gh_token:
+        print("Error: GH_TOKEN environment variable is not set.", file=sys.stderr)
+        sys.exit(1)
+
+    global HEADERS
+    HEADERS = {
+        "Authorization": f"Bearer {gh_token}",
+        "Accept": "application/vnd.github+json",
+        "X-GitHub-Api-Version": "2022-11-28",
+    }
+
+    collab_url = f"https://api.github.com/repos/{owner}/{name}/collaborators"
+    print(f"Fetching collaborators for {owner}/{name}...", file=sys.stderr)
+    collaborators = paginate_list(
+        collab_url, params={"per_page": 100, "affiliation": "all"}
+    )
+
+    rows: list[dict[str, Any]] = []
+    elevated = [c for c in collaborators if isinstance(c, dict) and has_write_plus(c)]
+    print(
+        f"Found {len(elevated)} collaborators with admin/maintain/write/triage.",
+        file=sys.stderr,
+    )
+
+    for i, col in enumerate(elevated, start=1):
+        login = col.get("login")
+        if not isinstance(login, str):
+            continue
+        print(f"  [{i}/{len(elevated)}] {login}", file=sys.stderr)
+
+        role = collaborator_role(col)
+        nickname = fetch_display_name(login)
+        cd = last_commit_date(owner, name, login)
+        issue_dt, pr_dt = last_issue_pr_dates(owner, name, login)
+        last_act_ymd = max_date_ymd(cd, issue_dt, pr_dt)
+        others = other_repos_activity_column(login, owner, name, days=args.events_days)
+        rows.append(
+            {
+                "_role_tier": role_sort_tier(col),
+                "github_username": login,
+                "nickname": nickname,
+                "role": role,
+                "last_activity_date": last_act_ymd,
+                "last_commit_date": iso_timestamp_to_ymd(cd),
+                "last_issue_date": iso_timestamp_to_ymd(issue_dt),
+                "last_pr_date": iso_timestamp_to_ymd(pr_dt),
+                "other_repos_90d": others,
+            }
+        )
+
+    def sort_key(r: dict[str, Any]) -> tuple[int, float]:
+        tier = r["_role_tier"]
+        act = parse_ymd(r.get("last_activity_date") or "")
+        ts = act.timestamp() if act else 0.0
+        return (tier, -ts)
+
+    rows.sort(key=sort_key)
+
+    fieldnames = [
+        "github_username",
+        "nickname",
+        "role",
+        "last_activity_date",
+        "last_commit_date",
+        "last_issue_date",
+        "last_pr_date",
+        "other_repos_90d",
+    ]
+    for r in rows:
+        del r["_role_tier"]
+    with open(args.output, "w", newline="", encoding="utf-8") as f:
+        w = csv.DictWriter(f, fieldnames=fieldnames)
+        w.writeheader()
+        w.writerows(rows)
+
+    print(f"Wrote {len(rows)} rows to {args.output}", file=sys.stderr)
+
+
+if __name__ == "__main__":
+    main()
--- a/third_party/sglang/.github/labeler.yml
+++ b/third_party/sglang/.github/labeler.yml
@@ -0,0 +1,122 @@
+# Configuration for the GitHub Labeler action
+# Automatically adds labels to PRs based on the files changed
+
+# Router specific (Rust code in sgl-model-gateway)
+model-gateway:
+  - changed-files:
+    - any-glob-to-any-file: 'sgl-model-gateway/**/*'
+
+# Kernel specific
+sgl-kernel:
+  - changed-files:
+    - any-glob-to-any-file: 'sgl-kernel/**/*'
+
+# JIT kernel specific
+jit-kernel:
+  - changed-files:
+    - any-glob-to-any-file: 'python/sglang/jit_kernel/**/*'
+
+# Documentation
+documentation:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*.md'
+      - 'docs/**/*'
+      - 'README*'
+
+# Dependencies
+dependencies:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/requirements*.txt'
+      - '**/Cargo.toml'
+      - '**/Cargo.lock'
+      - '**/pyproject*.toml'
+      - '**/setup.py'
+      - '**/poetry.lock'
+      - '**/package.json'
+      - '**/package-lock.json'
+
+# Multi-modal
+Multi-modal:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*multimodal*'
+      - '**/*vision*'
+      - '**/*vlm*'
+
+# Diffusion
+diffusion:
+  - changed-files:
+    - any-glob-to-any-file: 'python/sglang/multimodal_gen/**/*'
+
+# LoRA
+lora:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*lora*'
+
+# Quantization
+quant:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*quant*'
+      - '**/*quantization*'
+
+# Speculative decoding
+speculative-decoding:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*speculative*'
+
+# AMD specific
+amd:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*amd*'
+      - '**/*rocm*'
+
+# NPU specific
+npu:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*npu*'
+      - '**/*ascend*'
+
+# Blackwell
+blackwell:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*nvfp4*'
+      - 'sgl-kernel/csrc/attention/cutlass_sm100_mla/**/*'
+      - 'python/sglang/srt/layers/attention/trtllm_mla_backend.py'
+      - 'python/sglang/srt/layers/attention/trtllm_mha_backend.py'
+
+# DeepSeek specific
+deepseek:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*deepseek*'
+
+# HiCache
+hicache:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*hicache*'
+
+# Deterministic
+deterministic:
+  - changed-files:
+    - any-glob-to-any-file: 'python/sglang/srt/batch_invariant_ops/**/*'
+
+# Piecewise CUDA Graph
+piecewise-cuda-graph:
+  - changed-files:
+    - any-glob-to-any-file: 'python/sglang/srt/compilation/**/*'
+
+# Moore Threads specific
+mthreads:
+  - changed-files:
+    - any-glob-to-any-file:
+      - '**/*mthreads*'
+      - '**/*musa*'
--- a/third_party/sglang/.github/linters/lychee-ci.toml
+++ b/third_party/sglang/.github/linters/lychee-ci.toml
@@ -0,0 +1,42 @@
+no_progress = true
+verbose = "warn"
+timeout = 20
+max_concurrency = 8
+retry_wait_time = 2
+max_retries = 2
+
+# CI should validate external links over the network.
+offline = false
+scheme = ["http", "https"]
+
+exclude_path = [
+  # Exclude generated Sphinx build artifacts.
+  # - "(\\./)?" allows both "docs/..." and "./docs/..."
+  # - "[/\\\\]" supports both slash styles in CI environments
+  "^(\\./)?docs[/\\\\]_build[/\\\\]",
+]
+
+exclude = [
+  # Local-only endpoints referenced in docs/examples.
+  # These are expected to be unreachable in GitHub-hosted CI.
+  "^https?://localhost(:[0-9]+)?(/|$)",
+  "^http://127\\.0\\.0\\.1(:[0-9]+)?(/|$)",
+  # Vendor pages that frequently block/deny CI user-agents (transient 403/anti-bot).
+  "^https://www\\.intel\\.com/content/www/us/en/ark/products/series/240391/intel-arc-b-series-graphics\\.html$",
+  "^https://www\\.intel\\.com/content/www/us/en/ark/products/series/242616/intel-arc-pro-b-series-graphics\\.html$",
+  "^https://www\\.intel\\.com/content/www/us/en/products/sku/241598/intel-arc-b580-graphics/specifications\\.html$",
+
+  # Non-routable bind address used in examples, never externally reachable.
+  "^http://0\\.0\\.0\\.0(/|$)",
+
+  # Large doc portals with anti-bot/rate-limit behavior in CI.
+  # We keep API docs references in content but do not fail CI on access policy.
+  "^https://platform\\.openai\\.com/docs/",
+  "^https://gamma\\.app/docs/Optimizing-RL-with-SGLang-y0kqgj877k34779$",
+  "^https://aflah02\\.substack\\.com/p/multi-node-llm-inference-with-sglang/?$",
+
+  # Known noisy image URLs used in notebook-rendered examples.
+  "^https://github\\.com/sgl-project/sglang/blob/main/examples/assets/example_image\\.png\\?raw=true$",
+  "^https://raw\\.githubusercontent\\.com/sgl-project/sglang/main/examples/assets/example_image\\.png/?$",
+  "^https://raw\\.githubusercontent\\.com/sgl-project/sglang/main/assets/logo\\.png/?$",
+]
--- a/third_party/sglang/.github/linters/lychee.toml
+++ b/third_party/sglang/.github/linters/lychee.toml
@@ -0,0 +1,18 @@
+# .github/linters/lychee.toml
+no_progress = true
+verbose = "warn"
+timeout = 20
+max_concurrency = 8
+
+offline = true
+
+# Ignore generated docs output; check source docs only.
+exclude_path = [
+  "^(\\./)?docs[/\\\\]_build[/\\\\]",
+]
+
+exclude = [
+  "^https?://localhost(:[0-9]+)?(/|$)",
+  "^http://127\\.0\\.0\\.1(:[0-9]+)?(/|$)",
+  "^http://0\\.0\\.0\\.0(/|$)",
+]
--- a/third_party/sglang/.github/pull_request_template.md
+++ b/third_party/sglang/.github/pull_request_template.md
@@ -0,0 +1,33 @@
+<!-- Thank you for your contribution! Please follow these guidelines to enhance your pull request. If anything is unclear, submit your PR and reach out to maintainers for assistance. Join our Slack community at https://slack.sglang.io to discuss further. -->
+
+## Motivation
+
+<!-- Describe the purpose and goals of this pull request. -->
+
+## Modifications
+
+<!-- Detail the changes made in this pull request. -->
+
+## Accuracy Tests
+
+<!-- If this pull request affects model outputs (e.g., changes to the kernel or model forward code), provide accuracy test results. -->
+
+## Speed Tests and Profiling
+
+<!-- If this pull request impacts inference speed, provide benchmarking and profiling results. -->
+
+## Checklist
+
+- [ ] Format your code according to the [Format code with pre-commit](https://docs.sglang.io/developer_guide/contribution_guide.html#format-code-with-pre-commit).
+- [ ] Add unit tests according to the [Run and add unit tests](https://docs.sglang.io/developer_guide/contribution_guide.html#run-and-add-unit-tests).
+- [ ] Update documentation according to [Write documentations](https://docs.sglang.io/developer_guide/contribution_guide.html#write-documentations).
+- [ ] Provide accuracy and speed benchmark results according to [Test the accuracy](https://docs.sglang.io/developer_guide/contribution_guide.html#test-the-accuracy) and [Benchmark the speed](https://docs.sglang.io/developer_guide/contribution_guide.html#benchmark-the-speed).
+- [ ] Follow the SGLang code style [guidance](https://docs.sglang.io/developer_guide/contribution_guide.html#code-style-guidance).
+
+## Review and Merge Process
+
+1. Ping Merge Oncalls to start the process. See the [PR Merge Process](https://github.com/sgl-project/sglang/blob/main/.github/MAINTAINER.md#pull-request-merge-process).
+2. Get approvals from [CODEOWNERS](https://github.com/sgl-project/sglang/blob/main/.github/CODEOWNERS) and other reviewers.
+3. Trigger CI tests with [comments](https://docs.sglang.io/developer_guide/contribution_guide.html#how-to-trigger-ci-tests) or contact authorized users to do so.
+   - Common commands include `/tag-and-rerun-ci`, `/tag-run-ci-label`, `/rerun-failed-ci`
+4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.
--- a/third_party/sglang/.github/update_ci_permission.py
+++ b/third_party/sglang/.github/update_ci_permission.py
@@ -0,0 +1,244 @@
+"""
+Update the CI permissions configuration file.
+
+This script updates the `CI_PERMISSIONS.json` file, which defines the CI permissions granted to each user.
+
+The format of `CI_PERMISSIONS.json` is as follows:
+
+{
+    "username1": {
+        "can_tag_run_ci_label": true,
+        "can_rerun_failed_ci": true,
+        "cooldown_interval_minutes": 0,
+        "reason": "top contributor"
+    },
+    "username2": {
+        "can_tag_run_ci_label": true,
+        "can_rerun_failed_ci": true,
+        "cooldown_interval_minutes": 60,
+        "reason": "custom override"
+    }
+}
+
+Permissions are assigned according to the following rules:
+
+1. Add the top 50 contributors from the last 120 days with full permissions, no cooldown, and the reason "top contributor".
+2. Load all users from the existing `CI_PERMISSIONS.json` file and update their entries as follows:
+   - If a user is already covered by rule 1, skip that user.
+   - If the old reason of a user is "top contributor" but they are not in the current top contributors list, change their configuration to:
+       {
+           "can_tag_run_ci_label": true,
+           "can_rerun_failed_ci": true,
+           "cooldown_interval_minutes": 60,
+           "reason": "custom override"
+       }
+    - For all other cases, preserve the original configuration unchanged.
+3. All other users receive no permissions and a 120-minute cooldown (they are omitted from the file).
+
+Usage:
+    export GH_TOKEN="your_github_token"
+    python3 update_ci_permission.py
+
+    # Sort-only mode (no network calls, no GH_TOKEN required)
+    python3 update_ci_permission.py --sort-only
+"""
+
+import argparse
+import json
+import os
+from collections import Counter
+from datetime import datetime, timedelta, timezone
+
+try:
+    import requests
+except ImportError:
+    requests = None  # Only needed for non-sort-only runs
+
+# Configuration
+REPO_OWNER = "sgl-project"
+REPO_NAME = "sglang"
+FILE_NAME = os.path.join(os.path.dirname(__file__), "CI_PERMISSIONS.json")
+HEADERS = {}
+
+
+def github_api_get(endpoint, params=None):
+    """Helper to make paginated GitHub API requests."""
+    if requests is None:
+        raise RuntimeError(
+            "The requests package is required. Install it or use --sort-only."
+        )
+    if not HEADERS:
+        raise RuntimeError(
+            "GitHub headers not initialized. Set GH_TOKEN or use --sort-only."
+        )
+
+    results = []
+    url = f"https://api.github.com/repos/{REPO_OWNER}/{REPO_NAME}/{endpoint}"
+
+    while url:
+        response = requests.get(url, headers=HEADERS, params=params)
+        if response.status_code != 200:
+            print(f"Error fetching {url}: {response.status_code} {response.text}")
+            # If we fail to fetch, strictly return what we have or empty to avoid crashing logic
+            break
+
+        data = response.json()
+        if isinstance(data, list):
+            results.extend(data)
+        else:
+            return data  # Non-list response (not paginated usually)
+
+        # Handle pagination
+        url = None
+        if "link" in response.headers:
+            links = response.headers["link"].split(", ")
+            for link in links:
+                if 'rel="next"' in link:
+                    url = link[link.find("<") + 1 : link.find(">")]
+                    params = None  # Params are included in the next link
+                    break
+    return results
+
+
+def get_write_access_users():
+    """Fetches users with push (write) or admin access."""
+    print("Fetching collaborators with write access...")
+    # Note: This endpoint usually requires admin rights on the token.
+    collaborators = github_api_get("collaborators", params={"per_page": 100})
+
+    writers = set()
+    for col in collaborators:
+        perms = col.get("permissions", {})
+        # Check for admin, maintain, or push rights
+        if perms.get("admin") or perms.get("maintain") or perms.get("push"):
+            writers.add(col["login"])
+
+    print(f"Found {len(writers)} users with write access.")
+    return writers
+
+
+def get_top_contributors(days, limit):
+    """Fetches top contributors based on commit count in the last N days."""
+    print(f"Fetching commits from the last {days} days...")
+    since_date = (datetime.now(timezone.utc) - timedelta(days=days)).isoformat()
+
+    # Fetch commits
+    commits = github_api_get("commits", params={"since": since_date, "per_page": 100})
+
+    author_counts = Counter()
+    for commit in commits:
+        # commit['author'] contains the GitHub user object (can be None if not linked)
+        if commit.get("author") and "login" in commit["author"]:
+            author_counts[commit["author"]["login"]] += 1
+
+    top_users = [user for user, _ in author_counts.most_common(limit)]
+    print(f"Found {len(top_users)} top contributors in the last {days} days.")
+    return set(top_users)
+
+
+def load_existing_permissions():
+    if os.path.exists(FILE_NAME):
+        try:
+            with open(FILE_NAME, "r") as f:
+                return json.load(f)
+        except json.JSONDecodeError:
+            print(f"Warning: {FILE_NAME} is invalid JSON. Starting fresh.")
+    return {}
+
+
+def sort_permissions_file():
+    """Sort the existing CI permissions file alphabetically and exit."""
+    if not os.path.exists(FILE_NAME):
+        print(f"{FILE_NAME} not found. Nothing to sort.")
+        return
+
+    old_permissions = load_existing_permissions()
+    sorted_permissions = dict(sorted(old_permissions.items()))
+
+    with open(FILE_NAME, "w") as f:
+        json.dump(sorted_permissions, f, indent=4)
+        f.write("\n")
+
+    print(f"Sorted {FILE_NAME}. Total users: {len(sorted_permissions)}")
+
+
+def main():
+    parser = argparse.ArgumentParser(description="Update or sort CI permissions.")
+    parser.add_argument(
+        "--sort-only",
+        action="store_true",
+        help="Only sort CI_PERMISSIONS.json alphabetically without fetching data.",
+    )
+    args = parser.parse_args()
+
+    if args.sort_only:
+        sort_permissions_file()
+        return
+
+    gh_token = os.getenv("GH_TOKEN")
+    if not gh_token:
+        raise ValueError("Error: GH_TOKEN environment variable is not set.")
+
+    global HEADERS
+    HEADERS = {
+        "Authorization": f"Bearer {gh_token}",
+        "Accept": "application/vnd.github+json",
+        "X-GitHub-Api-Version": "2022-11-28",
+    }
+
+    # Gather Data
+    try:
+        write_access_users = get_write_access_users()
+    except Exception as e:
+        print(f"Warning: Could not fetch collaborators (check token scope). Error: {e}")
+        write_access_users = set()
+
+    top_contributors = get_top_contributors(days=120, limit=50)
+    old_permissions = load_existing_permissions()
+
+    new_permissions = {}
+
+    # Rule 1: Add Top 50 Contributors
+    for user in top_contributors:
+        new_permissions[user] = {
+            "can_tag_run_ci_label": True,
+            "can_rerun_failed_ci": True,
+            "can_rerun_stage": True,
+            "cooldown_interval_minutes": 0,
+            "reason": "top contributor",
+        }
+
+    # Rule 2: Process Existing Users (Merge Logic)
+    for user, config in old_permissions.items():
+        if user in new_permissions:
+            # Already handled by Rule 1 or 2
+            continue
+
+        old_reason = config.get("reason", "")
+
+        # If they fell off the top contributor list
+        if old_reason in ["top contributor"]:
+            new_permissions[user] = {
+                "can_tag_run_ci_label": True,
+                "can_rerun_failed_ci": True,
+                "can_rerun_stage": True,
+                "cooldown_interval_minutes": 60,
+                "reason": "custom override",
+            }
+        else:
+            # Preserve custom overrides
+            new_permissions[user] = config
+
+    # Save and Sort
+    # Sorting keys for cleaner diffs
+    sorted_permissions = dict(sorted(new_permissions.items()))
+
+    with open(FILE_NAME, "w") as f:
+        json.dump(sorted_permissions, f, indent=4)
+        f.write("\n")  # Add trailing newline
+
+    print(f"Successfully updated {FILE_NAME}. Total users: {len(sorted_permissions)}")
+
+
+if __name__ == "__main__":
+    main()
--- a/third_party/sglang/.github/workflows/amd-aiter-scout.yml
+++ b/third_party/sglang/.github/workflows/amd-aiter-scout.yml
@@ -0,0 +1,161 @@
+name: AMD AITER Scout
+
+on:
+  schedule:
+    - cron: '0 20 * * 1'   # Monday 20:00 UTC
+    - cron: '0 20 * * 4'   # Thursday 20:00 UTC
+  workflow_dispatch:
+    inputs:
+      aiter_ref:
+        description: 'AITER git ref (branch, tag, or SHA). Default: main (latest commit)'
+        required: false
+        type: string
+        default: 'main'
+      job_filter:
+        description: 'Comma-separated workflows to run: nightly-amd, nightly-amd-rocm720, pr-test-amd, pr-test-amd-rocm720. Default: all'
+        required: false
+        type: string
+        default: 'all'
+      continue_on_error:
+        description: 'Continue running other workflows even if one fails'
+        required: false
+        type: boolean
+        default: true
+
+concurrency:
+  group: amd-aiter-scout-${{ github.run_id }}
+  cancel-in-progress: true
+
+jobs:
+  resolve-aiter:
+    runs-on: ubuntu-latest
+    outputs:
+      aiter_sha: ${{ steps.resolve.outputs.sha }}
+      run_nightly_amd: ${{ steps.parse.outputs.run_nightly_amd }}
+      run_nightly_amd_rocm720: ${{ steps.parse.outputs.run_nightly_amd_rocm720 }}
+      run_pr_test_amd: ${{ steps.parse.outputs.run_pr_test_amd }}
+      run_pr_test_amd_rocm720: ${{ steps.parse.outputs.run_pr_test_amd_rocm720 }}
+    steps:
+      - name: Resolve AITER commit
+        id: resolve
+        run: |
+          REF="${{ inputs.aiter_ref || 'main' }}"
+          echo "Resolving AITER ref: ${REF}"
+
+          SHA=$(git ls-remote https://github.com/ROCm/aiter.git "refs/heads/${REF}" | head -1 | cut -f1)
+          if [ -z "$SHA" ]; then
+            SHA=$(git ls-remote https://github.com/ROCm/aiter.git "refs/tags/${REF}" | head -1 | cut -f1)
+          fi
+          if [ -z "$SHA" ]; then
+            SHA=$(git ls-remote https://github.com/ROCm/aiter.git "${REF}" | head -1 | cut -f1)
+          fi
+          if [ -z "$SHA" ]; then
+            SHA="${REF}"
+          fi
+
+          echo "sha=${SHA}" >> $GITHUB_OUTPUT
+          echo "### AITER Ref Resolution" >> $GITHUB_STEP_SUMMARY
+          echo "- **Requested ref:** \`${REF}\`" >> $GITHUB_STEP_SUMMARY
+          echo "- **Resolved SHA:** \`${SHA}\`" >> $GITHUB_STEP_SUMMARY
+          echo "- **AITER commit:** https://github.com/ROCm/aiter/commit/${SHA}" >> $GITHUB_STEP_SUMMARY
+
+      - name: Parse job filter
+        id: parse
+        run: |
+          FILTER="${{ inputs.job_filter || 'all' }}"
+          echo "Job filter: ${FILTER}"
+
+          if [[ "$FILTER" == "all" ]]; then
+            echo "run_nightly_amd=true" >> $GITHUB_OUTPUT
+            echo "run_nightly_amd_rocm720=true" >> $GITHUB_OUTPUT
+            echo "run_pr_test_amd=true" >> $GITHUB_OUTPUT
+            echo "run_pr_test_amd_rocm720=true" >> $GITHUB_OUTPUT
+          else
+            # Wrap with commas for exact substring matching (avoids "nightly-amd" matching "nightly-amd-rocm720")
+            PADDED=",${FILTER// /},"
+            echo "run_nightly_amd=$(echo "$PADDED" | grep -q ',nightly-amd,' && echo true || echo false)" >> $GITHUB_OUTPUT
+            echo "run_nightly_amd_rocm720=$(echo "$PADDED" | grep -q ',nightly-amd-rocm720,' && echo true || echo false)" >> $GITHUB_OUTPUT
+            echo "run_pr_test_amd=$(echo "$PADDED" | grep -q ',pr-test-amd,' && echo true || echo false)" >> $GITHUB_OUTPUT
+            echo "run_pr_test_amd_rocm720=$(echo "$PADDED" | grep -q ',pr-test-amd-rocm720,' && echo true || echo false)" >> $GITHUB_OUTPUT
+          fi
+
+          echo "### Job Filter" >> $GITHUB_STEP_SUMMARY
+          echo "- **Filter:** \`${FILTER}\`" >> $GITHUB_STEP_SUMMARY
+
+  call-nightly-amd:
+    if: needs.resolve-aiter.outputs.run_nightly_amd == 'true'
+    needs: resolve-aiter
+    uses: ./.github/workflows/nightly-test-amd.yml
+    secrets: inherit
+    with:
+      ref: ${{ github.sha }}
+      aiter_ref: ${{ needs.resolve-aiter.outputs.aiter_sha }}
+      job_filter: 'all'
+      continue_on_error: ${{ inputs.continue_on_error == '' && true || inputs.continue_on_error }}
+
+  call-nightly-amd-rocm720:
+    if: needs.resolve-aiter.outputs.run_nightly_amd_rocm720 == 'true'
+    needs: resolve-aiter
+    uses: ./.github/workflows/nightly-test-amd-rocm720.yml
+    secrets: inherit
+    with:
+      ref: ${{ github.sha }}
+      aiter_ref: ${{ needs.resolve-aiter.outputs.aiter_sha }}
+      job_filter: 'all'
+      continue_on_error: ${{ inputs.continue_on_error == '' && true || inputs.continue_on_error }}
+
+  call-pr-test-amd:
+    if: needs.resolve-aiter.outputs.run_pr_test_amd == 'true'
+    needs: resolve-aiter
+    uses: ./.github/workflows/pr-test-amd.yml
+    secrets: inherit
+    with:
+      run_all_tests: true
+      aiter_ref: ${{ needs.resolve-aiter.outputs.aiter_sha }}
+      continue_on_error: ${{ inputs.continue_on_error == '' && true || inputs.continue_on_error }}
+
+  call-pr-test-amd-rocm720:
+    if: needs.resolve-aiter.outputs.run_pr_test_amd_rocm720 == 'true'
+    needs: resolve-aiter
+    uses: ./.github/workflows/pr-test-amd-rocm720.yml
+    secrets: inherit
+    with:
+      run_all_tests: true
+      aiter_ref: ${{ needs.resolve-aiter.outputs.aiter_sha }}
+      continue_on_error: ${{ inputs.continue_on_error == '' && true || inputs.continue_on_error }}
+
+  check-all-jobs:
+    if: always()
+    needs:
+      - resolve-aiter
+      - call-nightly-amd
+      - call-nightly-amd-rocm720
+      - call-pr-test-amd
+      - call-pr-test-amd-rocm720
+    runs-on: ubuntu-latest
+    steps:
+      - name: Summary
+        run: |
+          echo "## AMD AITER Scout Results" >> $GITHUB_STEP_SUMMARY
+          echo "" >> $GITHUB_STEP_SUMMARY
+          echo "- **AITER SHA:** \`${{ needs.resolve-aiter.outputs.aiter_sha }}\`" >> $GITHUB_STEP_SUMMARY
+          echo "- **AITER commit:** https://github.com/ROCm/aiter/commit/${{ needs.resolve-aiter.outputs.aiter_sha }}" >> $GITHUB_STEP_SUMMARY
+          echo "" >> $GITHUB_STEP_SUMMARY
+          echo "| Workflow | Result |" >> $GITHUB_STEP_SUMMARY
+          echo "|----------|--------|" >> $GITHUB_STEP_SUMMARY
+          echo "| Nightly AMD (AITER Latest) | \`${{ needs.call-nightly-amd.result }}\` |" >> $GITHUB_STEP_SUMMARY
+          echo "| Nightly AMD ROCm 7.2 | \`${{ needs.call-nightly-amd-rocm720.result }}\` |" >> $GITHUB_STEP_SUMMARY
+          echo "| PR Test AMD (AITER Latest) | \`${{ needs.call-pr-test-amd.result }}\` |" >> $GITHUB_STEP_SUMMARY
+          echo "| PR Test AMD ROCm 7.2 | \`${{ needs.call-pr-test-amd-rocm720.result }}\` |" >> $GITHUB_STEP_SUMMARY
+
+      - name: Check if any job failed
+        run: |
+          if [[ "${{ contains(needs.*.result, 'failure') }}" == "true" ]]; then
+            echo "One or more workflows failed"
+            exit 1
+          fi
+          if [[ "${{ contains(needs.*.result, 'cancelled') }}" == "true" ]]; then
+            echo "One or more workflows were cancelled"
+            exit 1
+          fi
+          echo "All workflows passed"
--- a/third_party/sglang/.github/workflows/amd-ci-job-monitor.yml
+++ b/third_party/sglang/.github/workflows/amd-ci-job-monitor.yml
@@ -0,0 +1,338 @@
+name: AMD CI Job Monitor
+
+on:
+  schedule:
+    - cron: '0 0 * * *'  # Daily at midnight UTC
+  pull_request:
+    paths:
+      - '.github/workflows/amd-ci-job-monitor.yml'
+      - 'scripts/ci/utils/query_job_status.py'
+  workflow_dispatch:
+    inputs:
+      hours:
+        description: 'Time window in hours'
+        required: false
+        default: '24'
+        type: string
+      job_filter:
+        description: 'Job name filter (leave empty for all AMD jobs)'
+        required: false
+        type: string
+
+jobs:
+  fetch-actions-data:
+    name: Fetch Actions Snapshot
+    runs-on: ubuntu-latest
+    env:
+      GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Select workflows for snapshot
+        id: select-workflows
+        run: |
+          if [[ -n "${{ inputs.job_filter }}" ]]; then
+            echo "workflows=pr-test-amd.yml" >> "$GITHUB_OUTPUT"
+          else
+            echo "workflows=pr-test-amd.yml,nightly-test-amd.yml,pr-test-amd-rocm720.yml,nightly-test-amd-rocm720.yml" >> "$GITHUB_OUTPUT"
+          fi
+
+      - name: Fetch Actions data snapshot
+        timeout-minutes: 30
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --workflow "${{ steps.select-workflows.outputs.workflows }}" \
+            --hours ${{ inputs.hours || '24' }} \
+            --dump-data-file actions-job-snapshot.json
+
+      - name: Upload Actions data snapshot
+        uses: actions/upload-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: actions-job-snapshot.json
+          if-no-files-found: error
+
+  # Single job filter mode
+  custom-report:
+    name: Custom Job Report
+    if: ${{ inputs.job_filter }}
+    needs: fetch-actions-data
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Download Actions data snapshot
+        uses: actions/download-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: ci-data
+
+      - name: Generate Custom Job Report
+        timeout-minutes: 30
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --job "${{ inputs.job_filter }}" \
+            --workflow "pr-test-amd.yml" \
+            --hours ${{ inputs.hours || '24' }} \
+            --input-data-file ci-data/actions-job-snapshot.json \
+            --summary
+
+  # Parse workflow files to get job names dynamically
+  parse-workflows:
+    name: Parse Workflow Jobs
+    if: ${{ !inputs.job_filter }}
+    runs-on: ubuntu-latest
+    outputs:
+      pr_jobs: ${{ steps.parse.outputs.pr_jobs }}
+      nightly_jobs: ${{ steps.parse.outputs.nightly_jobs }}
+      pr_rocm720_jobs: ${{ steps.parse.outputs.pr_rocm720_jobs }}
+      nightly_rocm720_jobs: ${{ steps.parse.outputs.nightly_rocm720_jobs }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Parse workflow files
+        id: parse
+        run: |
+          # Parse pr-test-amd.yml and extract job names (exclude utility jobs)
+          # Excluded: call-gate, check-changes, pr-test-amd-finish, cancel, check-all-jobs
+          pr_jobs=$(yq -r '.jobs | keys | .[]' .github/workflows/pr-test-amd.yml | \
+            grep -v -E '^(call-gate|check-changes|pr-test-amd-finish|cancel|check-all-jobs)$' | \
+            jq -R -s -c 'split("\n") | map(select(length > 0))')
+          echo "pr_jobs=$pr_jobs" >> $GITHUB_OUTPUT
+          echo "PR jobs: $pr_jobs"
+
+          # Parse nightly-test-amd.yml and extract job names (exclude utility jobs)
+          # Excluded: check-all-jobs
+          nightly_jobs=$(yq -r '.jobs | keys | .[]' .github/workflows/nightly-test-amd.yml | \
+            grep -v -E '^(check-all-jobs)$' | \
+            jq -R -s -c 'split("\n") | map(select(length > 0))')
+          echo "nightly_jobs=$nightly_jobs" >> $GITHUB_OUTPUT
+          echo "Nightly jobs: $nightly_jobs"
+
+          # Parse pr-test-amd-rocm720.yml (exclude utility jobs)
+          # Excluded: call-gate, check-changes, pr-test-amd-finish, cancel, check-all-jobs
+          pr_rocm720_jobs=$(yq -r '.jobs | keys | .[]' .github/workflows/pr-test-amd-rocm720.yml | \
+            grep -v -E '^(call-gate|check-changes|pr-test-amd-finish|cancel|check-all-jobs)$' | \
+            jq -R -s -c 'split("\n") | map(select(length > 0))')
+          echo "pr_rocm720_jobs=$pr_rocm720_jobs" >> $GITHUB_OUTPUT
+          echo "PR ROCm 7.2 jobs: $pr_rocm720_jobs"
+
+          # Parse nightly-test-amd-rocm720.yml (exclude utility jobs)
+          # Excluded: check-all-jobs
+          nightly_rocm720_jobs=$(yq -r '.jobs | keys | .[]' .github/workflows/nightly-test-amd-rocm720.yml | \
+            grep -v -E '^(check-all-jobs)$' | \
+            jq -R -s -c 'split("\n") | map(select(length > 0))')
+          echo "nightly_rocm720_jobs=$nightly_rocm720_jobs" >> $GITHUB_OUTPUT
+          echo "Nightly ROCm 7.2 jobs: $nightly_rocm720_jobs"
+
+  # PR CI reports using dynamic matrix
+  pr-ci-reports:
+    name: PR - ${{ matrix.job_name }}
+    needs: [parse-workflows, fetch-actions-data]
+    if: ${{ !inputs.job_filter }}
+    runs-on: ubuntu-latest
+    strategy:
+      fail-fast: false
+      matrix:
+        job_name: ${{ fromJson(needs.parse-workflows.outputs.pr_jobs) }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Download Actions data snapshot
+        uses: actions/download-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: ci-data
+
+      - name: Generate Report
+        timeout-minutes: 15
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --job "${{ matrix.job_name }}" \
+            --workflow "pr-test-amd.yml" \
+            --hours ${{ inputs.hours || '24' }} \
+            --input-data-file ci-data/actions-job-snapshot.json \
+            --summary
+
+  # Nightly AMD test reports using dynamic matrix
+  nightly-reports:
+    name: Nightly - ${{ matrix.job_name }}
+    needs: [parse-workflows, fetch-actions-data]
+    if: ${{ !inputs.job_filter }}
+    runs-on: ubuntu-latest
+    strategy:
+      fail-fast: false
+      matrix:
+        job_name: ${{ fromJson(needs.parse-workflows.outputs.nightly_jobs) }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Download Actions data snapshot
+        uses: actions/download-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: ci-data
+
+      - name: Generate Nightly Report
+        timeout-minutes: 15
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --job "${{ matrix.job_name }}" \
+            --workflow "nightly-test-amd.yml" \
+            --hours ${{ inputs.hours || '24' }} \
+            --input-data-file ci-data/actions-job-snapshot.json \
+            --summary
+
+  # PR ROCm 7.2 CI reports using dynamic matrix
+  pr-rocm720-ci-reports:
+    name: PR ROCm720 - ${{ matrix.job_name }}
+    needs: [parse-workflows, fetch-actions-data]
+    if: ${{ !inputs.job_filter }}
+    runs-on: ubuntu-latest
+    strategy:
+      fail-fast: false
+      matrix:
+        job_name: ${{ fromJson(needs.parse-workflows.outputs.pr_rocm720_jobs) }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Download Actions data snapshot
+        uses: actions/download-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: ci-data
+
+      - name: Generate PR ROCm 7.2 Report
+        timeout-minutes: 15
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --job "${{ matrix.job_name }}" \
+            --workflow "pr-test-amd-rocm720.yml" \
+            --hours ${{ inputs.hours || '24' }} \
+            --input-data-file ci-data/actions-job-snapshot.json \
+            --summary
+
+  # Nightly ROCm 7.2 reports using dynamic matrix
+  nightly-rocm720-reports:
+    name: Nightly ROCm720 - ${{ matrix.job_name }}
+    needs: [parse-workflows, fetch-actions-data]
+    if: ${{ !inputs.job_filter }}
+    runs-on: ubuntu-latest
+    strategy:
+      fail-fast: false
+      matrix:
+        job_name: ${{ fromJson(needs.parse-workflows.outputs.nightly_rocm720_jobs) }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Download Actions data snapshot
+        uses: actions/download-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: ci-data
+
+      - name: Generate Nightly ROCm 7.2 Report
+        timeout-minutes: 15
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --job "${{ matrix.job_name }}" \
+            --workflow "nightly-test-amd-rocm720.yml" \
+            --hours ${{ inputs.hours || '24' }} \
+            --input-data-file ci-data/actions-job-snapshot.json \
+            --summary
+
+  # Runner fleet report - cross-workflow runner analytics in a single pass
+  runner-fleet-report:
+    name: Runner Fleet Report
+    if: ${{ !inputs.job_filter }}
+    needs: fetch-actions-data
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install dependencies
+        run: pip install tabulate
+
+      - name: Download Actions data snapshot
+        uses: actions/download-artifact@v4
+        with:
+          name: actions-job-snapshot
+          path: ci-data
+
+      - name: Generate Runner Fleet Report
+        timeout-minutes: 30
+        run: |
+          python scripts/ci/utils/query_job_status.py \
+            --repo ${{ github.repository }} \
+            --runner-report \
+            --workflow "pr-test-amd.yml,nightly-test-amd.yml,pr-test-amd-rocm720.yml,nightly-test-amd-rocm720.yml" \
+            --hours ${{ inputs.hours || '24' }} \
+            --input-data-file ci-data/actions-job-snapshot.json \
+            --summary
--- a/third_party/sglang/.github/workflows/auto-tune.yml
+++ b/third_party/sglang/.github/workflows/auto-tune.yml
@@ -0,0 +1,10 @@
+name: Auto tune
+
+on:
+  workflow_dispatch:
+
+jobs:
+  auto-tune-lint:
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
--- a/third_party/sglang/.github/workflows/bot-bump-flashinfer-version.yml
+++ b/third_party/sglang/.github/workflows/bot-bump-flashinfer-version.yml
@@ -0,0 +1,50 @@
+name: Bot Bump Flashinfer Version
+
+on:
+  workflow_dispatch:
+    inputs:
+      new_version:
+        description: 'New flashinfer version (e.g., 0.6.4)'
+        required: true
+        type: string
+
+permissions:
+  contents: write
+  pull-requests: write
+
+jobs:
+  bump-flashinfer-version:
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+        with:
+          token: ${{ secrets.GITHUB_TOKEN }}
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install Python dependencies
+        run: |
+            pip install tomli
+
+      - name: Configure Git and branch
+        run: |
+          git config user.name "sglang-bot"
+          git config user.email "sglang-bot@users.noreply.github.com"
+          RANDOM_SUFFIX=$(echo $RANDOM | md5sum | head -c 4)
+          BRANCH_NAME="bot/bump-flashinfer-version-${{ github.event.inputs.new_version }}-${RANDOM_SUFFIX}"
+          git checkout -b "$BRANCH_NAME"
+          echo "BRANCH_NAME=$BRANCH_NAME" >> $GITHUB_ENV
+
+      - name: Run flashinfer version bump script
+        run: |
+          python scripts/release/bump_flashinfer_version.py "${{ github.event.inputs.new_version }}"
+
+      - name: Commit and create PR
+        env:
+          GH_TOKEN: ${{ secrets.GH_PAT_FOR_PULL_REQUEST }}
+        run: |
+          bash scripts/release/commit_and_pr.sh "flashinfer" "${{ github.event.inputs.new_version }}" "$BRANCH_NAME"
--- a/third_party/sglang/.github/workflows/bot-bump-kernel-version-to-sglang.yml
+++ b/third_party/sglang/.github/workflows/bot-bump-kernel-version-to-sglang.yml
@@ -0,0 +1,60 @@
+name: Bot Bump Kernel Version to SGLang
+
+on:
+  workflow_dispatch:
+
+permissions:
+  contents: write
+  pull-requests: write
+
+jobs:
+  bump-kernel-version-to-sglang:
+    runs-on: ubuntu-latest
+    outputs:
+      branch_name: ${{ steps.set_output.outputs.branch_name }}
+      needs_sync: ${{ steps.check_sync.outputs.needs_sync }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+        with:
+          token: ${{ secrets.GITHUB_TOKEN }}
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install Python dependencies
+        run: |
+          pip install tomli
+
+      - name: Check if sync is needed
+        id: check_sync
+        run: |
+          python scripts/release/check_kernel_version_to_sglang.py
+
+      - name: Configure Git and branch
+        if: steps.check_sync.outputs.needs_sync == 'true'
+        id: set_output
+        run: |
+          git config user.name "sglang-bot"
+          git config user.email "sglang-bot@users.noreply.github.com"
+          RANDOM_SUFFIX=$(echo $RANDOM | md5sum | head -c 4)
+          KERNEL_VERSION="${{ steps.check_sync.outputs.kernel_version }}"
+          BRANCH_NAME="bot/bump-kernel-version-to-sglang-${KERNEL_VERSION}-${RANDOM_SUFFIX}"
+          git checkout -b "$BRANCH_NAME"
+          echo "BRANCH_NAME=$BRANCH_NAME" >> $GITHUB_ENV
+          echo "KERNEL_VERSION=$KERNEL_VERSION" >> $GITHUB_ENV
+          echo "branch_name=$BRANCH_NAME" >> $GITHUB_OUTPUT
+
+      - name: Run kernel version bump script
+        if: steps.check_sync.outputs.needs_sync == 'true'
+        run: |
+          python scripts/release/bump_kernel_version_to_sglang.py
+
+      - name: Commit and create PR
+        if: steps.check_sync.outputs.needs_sync == 'true'
+        env:
+          GH_TOKEN: ${{ secrets.GH_PAT_FOR_PULL_REQUEST }}
+        run: |
+          bash scripts/release/commit_and_pr_kernel_to_sglang.sh "$KERNEL_VERSION" "$BRANCH_NAME"
--- a/third_party/sglang/.github/workflows/bot-bump-kernel-version.yml
+++ b/third_party/sglang/.github/workflows/bot-bump-kernel-version.yml
@@ -0,0 +1,50 @@
+name: Bot Bump Kernel Version
+
+on:
+  workflow_dispatch:
+    inputs:
+      new_version:
+        description: 'New sgl-kernel version (e.g., 0.3.12)'
+        required: true
+        type: string
+
+permissions:
+  contents: write
+  pull-requests: write
+
+jobs:
+  bump-kernel-version:
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+        with:
+          token: ${{ secrets.GITHUB_TOKEN }}
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install Python dependencies
+        run: |
+          pip install tomli
+
+      - name: Configure Git and branch
+        run: |
+          git config user.name "sglang-bot"
+          git config user.email "sglang-bot@users.noreply.github.com"
+          RANDOM_SUFFIX=$(echo $RANDOM | md5sum | head -c 4)
+          BRANCH_NAME="bot/bump-kernel-version-${{ github.event.inputs.new_version }}-${RANDOM_SUFFIX}"
+          git checkout -b "$BRANCH_NAME"
+          echo "BRANCH_NAME=$BRANCH_NAME" >> $GITHUB_ENV
+
+      - name: Run kernel version bump script
+        run: |
+          python scripts/release/bump_kernel_version.py "${{ github.event.inputs.new_version }}"
+
+      - name: Commit and create PR
+        env:
+          GH_TOKEN: ${{ secrets.GH_PAT_FOR_PULL_REQUEST }}
+        run: |
+          bash scripts/release/commit_and_pr.sh "sgl-kernel" "${{ github.event.inputs.new_version }}" "$BRANCH_NAME"
--- a/third_party/sglang/.github/workflows/bot-bump-sglang-version.yml
+++ b/third_party/sglang/.github/workflows/bot-bump-sglang-version.yml
@@ -0,0 +1,89 @@
+name: Bot Bump SGLang Version
+
+on:
+  workflow_dispatch:
+    inputs:
+      new_version:
+        description: 'New SGLang version (e.g., 0.5.3 or 0.5.3rc0)'
+        required: true
+        type: string
+
+permissions:
+  contents: write
+  pull-requests: write
+
+jobs:
+  bump-sglang-version:
+    runs-on: ubuntu-latest
+    outputs:
+      branch_name: ${{ steps.set_output.outputs.branch_name }}
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+        with:
+          token: ${{ secrets.GITHUB_TOKEN }}
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Install Python dependencies
+        run: |
+          pip install tomli
+
+      - name: Configure Git and branch
+        id: set_output
+        run: |
+          git config user.name "sglang-bot"
+          git config user.email "sglang-bot@users.noreply.github.com"
+          RANDOM_SUFFIX=$(echo $RANDOM | md5sum | head -c 4)
+          BRANCH_NAME="bot/bump-sglang-version-${{ github.event.inputs.new_version }}-${RANDOM_SUFFIX}"
+          git checkout -b "$BRANCH_NAME"
+          echo "BRANCH_NAME=$BRANCH_NAME" >> $GITHUB_ENV
+          echo "branch_name=$BRANCH_NAME" >> $GITHUB_OUTPUT
+
+      - name: Run SGLang version bump script
+        run: |
+          python scripts/release/bump_sglang_version.py "${{ github.event.inputs.new_version }}"
+
+      - name: Commit and create PR
+        env:
+          GH_TOKEN: ${{ secrets.GH_PAT_FOR_PULL_REQUEST }}
+        run: |
+          bash scripts/release/commit_and_pr.sh "SGLang" "${{ github.event.inputs.new_version }}" "$BRANCH_NAME"
+
+  run-nightly-tests-nvidia:
+    needs: bump-sglang-version
+    uses: ./.github/workflows/nightly-test-nvidia.yml
+    with:
+      ref: ${{ needs.bump-sglang-version.outputs.branch_name }}
+    secrets: inherit
+
+  run-nightly-tests-amd:
+    needs: bump-sglang-version
+    uses: ./.github/workflows/nightly-test-amd.yml
+    with:
+      ref: ${{ needs.bump-sglang-version.outputs.branch_name }}
+    secrets: inherit
+
+  run-nightly-tests-npu:
+    needs: bump-sglang-version
+    uses: ./.github/workflows/nightly-test-npu.yml
+    with:
+      ref: ${{ needs.bump-sglang-version.outputs.branch_name }}
+    secrets: inherit
+
+  run-pr-tests-xeon:
+    needs: bump-sglang-version
+    uses: ./.github/workflows/pr-test-xeon.yml
+    with:
+      ref: ${{ needs.bump-sglang-version.outputs.branch_name }}
+    secrets: inherit
+
+  run-pr-tests-xpu:
+    needs: bump-sglang-version
+    uses: ./.github/workflows/pr-test-xpu.yml
+    with:
+      ref: ${{ needs.bump-sglang-version.outputs.branch_name }}
+    secrets: inherit
--- a/third_party/sglang/.github/workflows/bot-cherry-pick.yml
+++ b/third_party/sglang/.github/workflows/bot-cherry-pick.yml
@@ -0,0 +1,182 @@
+name: Bot Cherry Pick to Release Branch
+
+on:
+  workflow_dispatch:
+    inputs:
+      commit_sha:
+        description: 'Commit SHA to cherry-pick (full or short hash)'
+        required: true
+        type: string
+      target_branch:
+        description: 'Target release branch (e.g., release/v0.5.7)'
+        required: true
+        type: string
+      create_pr:
+        description: 'Create a PR instead of pushing directly'
+        required: false
+        type: boolean
+        default: true
+
+permissions:
+  contents: write
+  pull-requests: write
+
+concurrency:
+  group: cherry-pick-${{ github.event.inputs.target_branch }}
+  cancel-in-progress: false
+
+jobs:
+  cherry-pick:
+    if: github.repository == 'sgl-project/sglang'
+    runs-on: ubuntu-latest
+    environment: 'prod'
+    steps:
+      - name: Validate inputs
+        env:
+          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
+        run: |
+          if [[ ! "$TARGET_BRANCH" =~ ^release/v[0-9]+\.[0-9]+(\.[0-9]+)?$ ]]; then
+            echo "::error::Target branch must match pattern 'release/vX.Y' or 'release/vX.Y.Z' (e.g., release/v0.5.7)"
+            exit 1
+          fi
+
+      - name: Checkout code
+        uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+          token: ${{ secrets.GH_PAT_FOR_PULL_REQUEST }}
+
+      - name: Configure Git
+        run: |
+          git config user.name "sglang-bot"
+          git config user.email "sglang-bot@users.noreply.github.com"
+
+      - name: Validate target branch exists
+        env:
+          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
+        run: |
+          git fetch origin
+          if ! git ls-remote --exit-code --heads origin "$TARGET_BRANCH" > /dev/null 2>&1; then
+            echo "::error::Target branch '$TARGET_BRANCH' does not exist on remote"
+            exit 1
+          fi
+
+      - name: Get commit info
+        id: commit_info
+        env:
+          COMMIT_SHA_INPUT: ${{ github.event.inputs.commit_sha }}
+        run: |
+          # Verify commit exists
+          if ! git cat-file -t "$COMMIT_SHA_INPUT" > /dev/null 2>&1; then
+            echo "::error::Commit SHA '$COMMIT_SHA_INPUT' does not exist"
+            exit 1
+          fi
+
+          # Get full SHA if short hash provided
+          FULL_SHA=$(git rev-parse "$COMMIT_SHA_INPUT")
+          COMMIT_TITLE=$(git log -1 --format="%s" "$FULL_SHA")
+          SHORT_SHA=$(git rev-parse --short "$FULL_SHA")
+          echo "full_sha=$FULL_SHA" >> $GITHUB_OUTPUT
+          echo "short_sha=$SHORT_SHA" >> $GITHUB_OUTPUT
+          # Use delimiter for multiline-safe output
+          {
+            echo "commit_title<<EOF"
+            echo "$COMMIT_TITLE"
+            echo "EOF"
+          } >> $GITHUB_OUTPUT
+          echo "Cherry-picking commit: $SHORT_SHA - $COMMIT_TITLE"
+
+      - name: Cherry-pick commit
+        id: cherry_pick
+        env:
+          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
+          FULL_SHA: ${{ steps.commit_info.outputs.full_sha }}
+          SHORT_SHA: ${{ steps.commit_info.outputs.short_sha }}
+          CREATE_PR: ${{ github.event.inputs.create_pr }}
+        run: |
+          if [[ "$CREATE_PR" == "true" ]]; then
+            # Create a new branch for the PR
+            RANDOM_SUFFIX=$(head -c 4 /dev/urandom | xxd -p)
+            NEW_BRANCH="cherry-pick/${SHORT_SHA}-to-${TARGET_BRANCH#release/}-${RANDOM_SUFFIX}"
+            git checkout -b "$NEW_BRANCH" "origin/$TARGET_BRANCH"
+            echo "new_branch=$NEW_BRANCH" >> $GITHUB_OUTPUT
+          else
+            # Checkout target branch directly
+            git checkout "$TARGET_BRANCH"
+          fi
+
+          # Attempt cherry-pick
+          if git cherry-pick "$FULL_SHA"; then
+            echo "cherry_pick_success=true" >> $GITHUB_OUTPUT
+          else
+            echo "::error::Cherry-pick failed due to conflicts. Please resolve manually."
+            git cherry-pick --abort || true
+            echo "cherry_pick_success=false" >> $GITHUB_OUTPUT
+            exit 1
+          fi
+
+      - name: Push changes
+        if: steps.cherry_pick.outputs.cherry_pick_success == 'true'
+        env:
+          CREATE_PR: ${{ github.event.inputs.create_pr }}
+          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
+          NEW_BRANCH: ${{ steps.cherry_pick.outputs.new_branch }}
+        run: |
+          if [[ "$CREATE_PR" == "true" ]]; then
+            git push origin "$NEW_BRANCH"
+          else
+            git push origin "$TARGET_BRANCH"
+          fi
+
+      - name: Create Pull Request
+        if: steps.cherry_pick.outputs.cherry_pick_success == 'true' && github.event.inputs.create_pr == 'true'
+        env:
+          GH_TOKEN: ${{ secrets.GH_PAT_FOR_PULL_REQUEST }}
+          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
+          SHORT_SHA: ${{ steps.commit_info.outputs.short_sha }}
+          COMMIT_TITLE: ${{ steps.commit_info.outputs.commit_title }}
+          FULL_SHA: ${{ steps.commit_info.outputs.full_sha }}
+          NEW_BRANCH: ${{ steps.cherry_pick.outputs.new_branch }}
+        run: |
+          PR_TITLE="[Cherry-pick] ${COMMIT_TITLE} to ${TARGET_BRANCH}"
+
+          gh pr create \
+            --title "$PR_TITLE" \
+            --base "$TARGET_BRANCH" \
+            --head "$NEW_BRANCH" \
+            --label "cherry-pick" \
+            --body-file - <<EOF
+          Cherry-pick of commit ${FULL_SHA} to \`${TARGET_BRANCH}\`
+
+          **Original commit:** ${FULL_SHA}
+          **Original title:** ${COMMIT_TITLE}
+
+          ---
+          *This PR was automatically created by the cherry-pick workflow.*
+          EOF
+
+      - name: Summary
+        if: always()
+        env:
+          FULL_SHA: ${{ steps.commit_info.outputs.full_sha }}
+          COMMIT_TITLE: ${{ steps.commit_info.outputs.commit_title }}
+          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
+          CHERRY_PICK_SUCCESS: ${{ steps.cherry_pick.outputs.cherry_pick_success }}
+          CREATE_PR: ${{ github.event.inputs.create_pr }}
+          NEW_BRANCH: ${{ steps.cherry_pick.outputs.new_branch }}
+          ACTOR: ${{ github.actor }}
+        run: |
+          echo "## Cherry-Pick Summary" >> $GITHUB_STEP_SUMMARY
+          echo "" >> $GITHUB_STEP_SUMMARY
+          echo "- **Triggered by:** @${ACTOR}" >> $GITHUB_STEP_SUMMARY
+          echo "- **Commit:** ${FULL_SHA}" >> $GITHUB_STEP_SUMMARY
+          echo "- **Title:** ${COMMIT_TITLE}" >> $GITHUB_STEP_SUMMARY
+          echo "- **Target Branch:** ${TARGET_BRANCH}" >> $GITHUB_STEP_SUMMARY
+          if [[ "$CHERRY_PICK_SUCCESS" == "true" ]]; then
+            echo "- **Status:** ✅ Success" >> $GITHUB_STEP_SUMMARY
+          else
+            echo "- **Status:** ❌ Failed" >> $GITHUB_STEP_SUMMARY
+          fi
+          if [[ "$CREATE_PR" == "true" && "$CHERRY_PICK_SUCCESS" == "true" ]]; then
+            echo "- **PR Branch:** ${NEW_BRANCH}" >> $GITHUB_STEP_SUMMARY
+          fi
--- a/third_party/sglang/.github/workflows/cancel-pr-workflow-on-merge.yml
+++ b/third_party/sglang/.github/workflows/cancel-pr-workflow-on-merge.yml
@@ -0,0 +1,22 @@
+name: Cancel PR Workflows on Merge
+
+on:
+  pull_request_target:
+    types:
+      - closed
+
+permissions:
+  actions: write
+
+jobs:
+  cancel:
+    if: github.event.pull_request.merged == true
+    runs-on: ubuntu-latest
+    steps:
+      - name: Cancel Previous Runs
+        uses: styfle/cancel-workflow-action@0.12.1
+        with:
+          workflow_id: all
+          access_token: ${{ secrets.GITHUB_TOKEN }}
+          ignore_sha: true
+          pr_number: ${{ github.event.pull_request.number }}
--- a/third_party/sglang/.github/workflows/cancel-unfinished-pr-tests.yml
+++ b/third_party/sglang/.github/workflows/cancel-unfinished-pr-tests.yml
@@ -0,0 +1,155 @@
+name: Cancel Unfinished PR Runs
+
+on:
+  workflow_dispatch:
+    inputs:
+      workflows:
+        description: 'Space-separated list of workflow filenames to cancel'
+        required: true
+        type: string
+        default: 'pr-test.yml'
+      include_high_priority:
+        description: 'Also cancel runs from high-priority PRs'
+        required: false
+        type: boolean
+        default: false
+
+permissions:
+  actions: write   # Needed to cancel runs
+  contents: read   # Needed to read repo info
+  pull-requests: read  # needed for gh pr view (labels)
+
+jobs:
+  cancel-unfinished-pr-runs:
+    runs-on: ubuntu-latest
+    steps:
+      - name: Install GitHub CLI
+        run: sudo apt-get install -y gh jq
+
+      - name: Cancel unfinished PR-associated runs (skip high-priority PRs)
+        env:
+          GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+          REPO: ${{ github.repository }}
+          WORKFLOWS: ${{ github.event.inputs.workflows || 'pr-test.yml' }}
+          INCLUDE_HIGH_PRIORITY: ${{ github.event.inputs.include_high_priority || 'false' }}
+        shell: bash
+        run: |
+          set -euo pipefail
+
+          # Read the space-separated string from the input into a bash array
+          read -r -a WORKFLOW_FILES <<< "${WORKFLOWS}"
+
+          echo "Targeting ${#WORKFLOW_FILES[@]} workflow(s): ${WORKFLOWS}"
+          echo ""
+
+          for workflow_file in "${WORKFLOW_FILES[@]}"; do
+            echo "========================================="
+            echo "Workflow: $workflow_file"
+            echo "========================================="
+
+            # Get all unfinished runs
+            all_runs=$(gh run list \
+              --repo "$REPO" \
+              --workflow "$workflow_file" \
+              --json databaseId,status,event,url,createdAt \
+              --limit 1000 \
+            | jq -c '.[] | select(.status=="queued" or .status=="waiting" or .status=="in_progress")')
+
+            if [ -z "$all_runs" ]; then
+              echo "✅ No unfinished runs found"
+              echo ""
+              continue
+            fi
+
+            # Count runs by event type
+            total_runs=$(echo "$all_runs" | wc -l)
+            pr_runs=$(echo "$all_runs" | jq -s '[.[] | select(.event=="pull_request")] | length')
+            other_runs=$(echo "$all_runs" | jq -s '[.[] | select(.event!="pull_request")] | length')
+
+            echo "📊 Summary: $total_runs unfinished runs ($pr_runs PR-related, $other_runs other)"
+            echo ""
+
+            # Process non-PR runs first
+            if [ "$other_runs" -gt 0 ]; then
+              echo "--- Non-PR Runs ---"
+              echo "$all_runs" | jq -c 'select(.event!="pull_request")' | while read -r run; do
+                run_url=$(echo "$run" | jq -r '.url')
+                run_event=$(echo "$run" | jq -r '.event')
+                run_status=$(echo "$run" | jq -r '.status')
+                echo "  • $run_event ($run_status): $run_url"
+              done
+              echo ""
+            fi
+
+            # Process PR runs
+            if [ "$pr_runs" -gt 0 ]; then
+              echo "--- PR Runs (checking for cancellation) ---"
+              echo "$all_runs" | jq -c 'select(.event=="pull_request")' | while read -r run; do
+                run_id=$(echo "$run" | jq -r '.databaseId')
+                run_url=$(echo "$run" | jq -r '.url')
+                run_status=$(echo "$run" | jq -r '.status')
+
+                echo ""
+                echo "Run ($run_status): $run_url"
+
+                # Fetch full run details to get head repository and branch info
+                run_details=$(gh api -H "Accept: application/vnd.github+json" \
+                  "repos/$REPO/actions/runs/$run_id" 2>/dev/null || true)
+
+                if [ -z "$run_details" ]; then
+                  echo "  ⚠️  Could not fetch run details, skipping"
+                  continue
+                fi
+
+                # Get head owner and branch (works for both fork and non-fork PRs)
+                head_owner=$(echo "$run_details" | jq -r '.head_repository.owner.login // empty')
+                head_branch=$(echo "$run_details" | jq -r '.head_branch // empty')
+
+                if [ -z "$head_owner" ] || [ -z "$head_branch" ]; then
+                  echo "  ⚠️  Missing head info, skipping"
+                  continue
+                fi
+
+                echo "  Branch: ${head_owner}:${head_branch}"
+
+                # Find PR by searching with head=owner:branch
+                pr_number=$(gh api -H "Accept: application/vnd.github+json" \
+                  "repos/$REPO/pulls?state=open&head=${head_owner}:${head_branch}" \
+                  --jq '.[0].number // empty' 2>/dev/null || true)
+
+                if [ -z "$pr_number" ]; then
+                  echo "  ⚠️  No open PR found, skipping"
+                  continue
+                fi
+
+                pr_url="https://github.com/$REPO/pull/$pr_number"
+                echo "  PR: $pr_url"
+
+                # Check for high priority label
+                labels=$(gh pr view "$pr_number" --repo "$REPO" --json labels \
+                  | jq -r '.labels[].name' 2>/dev/null || true)
+
+                if echo "$labels" | grep -Fxq "bypass-maintenance"; then
+                  echo "  🛑 Skipping (bypass-maintenance label, never cancelled)"
+                  continue
+                fi
+
+                if echo "$labels" | grep -Fxq "high priority"; then
+                  if [ "$INCLUDE_HIGH_PRIORITY" != "true" ]; then
+                    echo "  🛑 Skipping (high priority label)"
+                    continue
+                  fi
+                  echo "  ⚠️  High priority PR, but include_high_priority is enabled"
+                fi
+
+                echo "  🚫 Cancelling..."
+                gh run cancel "$run_id" --repo "$REPO" || echo "  ⚠️  Cancellation failed"
+              done
+            fi
+
+            echo ""
+          done
+
+          echo "========================================="
+          echo "✅ Processing complete"
+          echo "========================================="
--- a/third_party/sglang/.github/workflows/ci-coverage-overview.yml
+++ b/third_party/sglang/.github/workflows/ci-coverage-overview.yml
@@ -0,0 +1,154 @@
+name: CI Coverage Overview
+
+on:
+  schedule:
+    - cron: '0 6 * * *'  # Daily at 6 AM UTC
+  pull_request:
+    paths:
+      - '.github/workflows/ci-coverage-overview.yml'
+      - 'scripts/ci/utils/ci_coverage_report.py'
+      - 'test/registered/**'
+  workflow_dispatch:
+    inputs:
+      output_format:
+        description: 'Output format'
+        required: false
+        default: 'markdown'
+        type: choice
+        options:
+          - markdown
+          - json
+
+jobs:
+  summary:
+    name: Summary
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Generate Summary Report
+        run: |
+          python scripts/ci/utils/ci_coverage_report.py --section summary
+
+  by-folder:
+    name: Tests by Folder
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Generate Tests by Folder Report
+        run: |
+          python scripts/ci/utils/ci_coverage_report.py --section by-folder
+
+  by-suite:
+    name: Tests by Suite
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Generate Tests by Suite Report
+        run: |
+          python scripts/ci/utils/ci_coverage_report.py --section by-suite
+
+  unit-test-coverage:
+    name: Unit Test Code Coverage
+    if: github.event_name != 'pull_request'
+    runs-on: 1-gpu-h100
+    timeout-minutes: 30
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Install dependencies
+        timeout-minutes: 10
+        run: |
+          pip install -e "python/[test]"
+
+      - name: Run unit tests with coverage
+        timeout-minutes: 10
+        run: |
+          pytest test/registered/unit/ \
+            --cov --cov-config=.coveragerc \
+            --cov-report=term-missing:skip-covered \
+            --continue-on-collection-errors \
+            -v | tee coverage_output.txt
+
+      - name: Write coverage to summary
+        if: always()
+        run: |
+          echo "## Unit Test Code Coverage" >> $GITHUB_STEP_SUMMARY
+          echo "" >> $GITHUB_STEP_SUMMARY
+          echo "**Commit:** \`${GITHUB_SHA::8}\` | **Branch:** \`${GITHUB_REF_NAME}\`" >> $GITHUB_STEP_SUMMARY
+          echo "" >> $GITHUB_STEP_SUMMARY
+
+          # Test result line (e.g., "== 42 passed, 1 failed in 23.5s ==")
+          echo '```' >> $GITHUB_STEP_SUMMARY
+          grep -E '^=+.*passed' coverage_output.txt >> $GITHUB_STEP_SUMMARY || true
+          echo "" >> $GITHUB_STEP_SUMMARY
+          # Coverage total
+          grep -E '^TOTAL ' coverage_output.txt >> $GITHUB_STEP_SUMMARY || true
+          echo '```' >> $GITHUB_STEP_SUMMARY
+
+          # Partially covered core modules (1-49%) — most actionable for contributors
+          # Only show modules with testable logic; skip configs, models, layers, etc.
+          LOW_COV=$(awk '/^python\/.*%/ {
+            for (i=1; i<=NF; i++) {
+              if ($i ~ /^[0-9]+%$/) {
+                pct = $i + 0
+                if (pct >= 1 && pct < 50) printf "%-80s %5s  %s\n", $1, $(i-2), $i
+                break
+              }
+            }
+          }' coverage_output.txt \
+            | grep -E '/(mem_cache|managers|sampling|parser|observability|function_call|entrypoints|speculative|multimodal|utils)/' \
+            | head -40 || true)
+          if [ -n "$LOW_COV" ]; then
+            echo "" >> $GITHUB_STEP_SUMMARY
+            echo "<details><summary>Core modules with coverage below 50% — good candidates for more unit tests</summary>" >> $GITHUB_STEP_SUMMARY
+            echo "" >> $GITHUB_STEP_SUMMARY
+            echo '```' >> $GITHUB_STEP_SUMMARY
+            echo "$LOW_COV" >> $GITHUB_STEP_SUMMARY
+            echo '```' >> $GITHUB_STEP_SUMMARY
+            echo "</details>" >> $GITHUB_STEP_SUMMARY
+          fi
+
+  json-export:
+    name: JSON Export
+    runs-on: ubuntu-latest
+    if: inputs.output_format == 'json'
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.10'
+
+      - name: Generate JSON Report
+        run: |
+          python scripts/ci/utils/ci_coverage_report.py --output-format json > ci_coverage.json
+
+      - name: Upload JSON artifact
+        uses: actions/upload-artifact@v4
+        with:
+          name: ci-coverage-report
+          path: ci_coverage.json
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Gahow Wang	7568e041ff	docs(kvc): record real Ali KVC experiment results	2026-05-12 05:28:06 +00:00
Gahow Wang	4e8f943875	feat(kvc): add real Ali replay workflow	2026-05-12 05:28:00 +00:00
kzlin	1d51704dad	docs(kvc): agentic-fit analysis, refactor plan, validation report Three new docs covering the structural-fit investigation: - AGENTIC_FIT_ANALYSIS_ZH.md: §1-§7 of structural design issues that surface KVC vs vanilla DP gap on real agentic workloads (SWE 50sess). Quantifies session pinning, LRU shortfall, P-side imbalance, time-scale distortion, etc., with code citations and N=3 rerun data. - REFACTOR_PLAN_ZH.md: KISS-edition refactor plan. After verifying the original "estimate inflation" and "resident_blocks aging" claims were not real bugs, scope shrinks to one code change (backpressure) plus a 4-run smoke sweep within an 8h budget. - STRUCTURAL_VALIDATION_REPORT_ZH.md: validates §1-§7 claims using existing v5 baseline rerun data + 8DP CA baseline. Each claim labeled fully-supported / indirect / retracted with the data source. Notes that backpressure E2E validation is pending GPU smoke run. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-05-06 21:30:11 +08:00
kzlin	7affb565b2	feat(kvc): add backpressure smoke sweep + analyzer (and v6 p1 profile script) scripts/sweep_backpressure_smoke.sh: 4-run smoke matrix (KVC baseline / KVC + backpressure / KVC + backpressure @ time-scale=1 / DP @ time-scale=1) designed to fit ~3-4h GPU budget. Validates §3 backpressure implementation and partially probes §7 time-scale distortion. scripts/analysis/analyze_backpressure_smoke.py: consumes the new structural/* jsonl files plus request-metrics; emits headline metrics, backpressure histograms, admission probe stats, and per-session pinning distribution. scripts/sweep_tp1_v6_p1_profile.sh: pre-existing v6 P1 profile sweep script (was untracked; included for completeness). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-05-06 21:29:56 +08:00
kzlin	c47adaf8e3	feat(kvc): honor admission backpressure hints + structural event logging Replay-side changes paired with the SGLang admission hint: - DecodeResidencyState gains pause_until_s; admission probe parses recommended_pause_ms and updates the per-D pause window. - _wait_for_decode_pause is invoked at request entry points (_invoke_router, _invoke_session_direct) so requests stall before hitting a saturated D instead of timing out via mooncake. - New CLI flags: --enable-backpressure (default off, baseline preserved), --backpressure-max-pause-s (cap on per-request sleep, default 2s). Structural instrumentation written under <run_dir>/structural/: - admission-events.jsonl: every admission probe (RTT, queue_depth, pause_ms, available_tokens, evicted_count) - backpressure-events.jsonl: every actual pause sleep - session-d-binding.jsonl: per-request policy decision Used to validate the structural claims documented separately. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-05-06 21:29:46 +08:00
kzlin	ca4b64c79a	feat(sglang): expose backpressure pause hint in admit_direct_append Add `recommended_pause_ms` field to DirectAppendAdmissionReqOutput so D can advise callers when its transfer queue is heavy or KV pool is near capacity. The hint is computed from transfer_queue_depth, retracted_queue_depth, and post-trim token_usage; thresholds are simple heuristics (>0.90 usage, >=8 queue depth, retracted>0). Default behavior is unchanged for callers that ignore the field. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-05-06 21:29:30 +08:00
kzlin	4978c0d0cd	profile(kvc): rewrite v5+profile report after critic audit + P0/P1 instrument Hostile audit of the original report flagged three load-bearing errors: 1. held_tokens semantic was inverted. session_held_tokens() at session_aware_cache.py:278-282 sums (kv_allocated_len - cache_protected_len) per slot, i.e. slot-private (NOT in radix tree). So "other = cap - held - avail" actually CONTAINS the radix-tree protected prefix cache (likely the single biggest component for shared agentic prefixes), not just running batch + in-flight as the original report claimed. 2. Admission-race causal hypothesis for the 415 EXP2+profile errors is contradicted by the data: 414/415 errors have kv_transfer_blocks > 0 — they passed admission and died downstream ("generate stream ended before producing any token", raised by the client when a 200 response had an empty stream). 3. Polling deconfound was too quickly dismissed. Mode counts shift ~1:1 (session-cap-fb -356 / kvcache-centric +406), and /server_info is not a passive read — it dispatches into the scheduler main loop and iterates every session slot. Plus: per-D error% confounded by sticky session affinity (only 18 unique sessions cause 415 errors, decode-3 had 0 errors only because no high-error session landed there); decile 10 "recovery" was an equal-time binning artifact (24.5% under equal-count); v5 vs v5+profile time gap was 21h not 6h; p50/p90 latency comparison is N=1. Rewritten report (docs/V5_PROFILE_INVESTIGATION_ZH.md) marks each correction with ⚠️ and demotes admission-race to one of four hypotheses (H1-H4). Action items split into P0 (verify, must do first) and P1 (instrument): P0 — scripts/sweep_tp1_v5_baseline_rerun_exp2.sh runs 3x v5 baseline EXP2 (no polling, identical config to the original v5 run) to test whether the 9-error baseline result is reproducible. If 3 runs give ~9 errors and profile gives 415, polling is the leading suspect. Currently running in background. P1 — scheduler.py:_compute_pool_breakdown_for_diagnostics adds a read-only "pool_breakdown" dict to /server_info covering: radix_evictable_tokens, radix_protected_tokens, slot_private_held_tokens, session_slot_count, running_batch_{reqs,kv_tokens}, transfer_queue_{reqs,tokens}, prealloc_queue_{reqs,tokens}, retracted_queue_{reqs,tokens}. With these, "unaccounted = cap - sum(known)" exposes true leakage. replay.py captures all fields into the per-tick row; analyzer prints the decomposition and gracefully handles old timeseries (prints "P1 instrument absent"). Mock-tested end-to-end. SGLang patch is read-only and does not affect admission/scheduling. Old v5+profile data still analyzes correctly. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 22:29:21 +08:00
kzlin	51f5386691	profile(kvc): add D KV pool timeseries poller + analyzer for v6 root-cause v5 dropped errors but pushed session-cap fallback to 46-51%. Before adding v6 mitigations we need to attribute that capacity loss to one of: (a) active sessions — real footprint (b) idle-evictable sessions — LRU not aggressive enough (c) prefill backup blocks / in-flight / fragmentation — release timing Without this it's all guessing. Plumb a 1Hz poller into replay that hits each P/D worker's /server_info, captures session_cache + memory_usage, and writes a per-worker time-series JSONL to <run_dir>/d-pool-timeseries.jsonl. Off by default (--pool-poll-interval-s 0); v5+profile sweep enables it at 1.0s. Per-tick HTTP cost is ~8 parallel /server_info calls — negligible relative to the 50min run. Analyzer (scripts/analysis/analyze_pool_timeseries.py) decomposes each D's capacity into active_held / idle_evictable / other (= cap-held-avail, the backup-blocks bucket) / free, and reports session residency churn across workers as a starvation/thrashing signal. Mock-tested poller end-to-end (cancellation clean, file flushed, sessions captured); analyzer validated against synthetic timeseries. Next: run scripts/sweep_tp1_v5_optD_profile.sh on hardware (~90min), then analyze results to pick a v6 direction. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 20:04:21 +08:00
kzlin	6572d7f3f4	docs: add v5 chapter (Option D worker-mode admission) and rename to V1_TO_V5 v5 sweep (sweep_tp1_v5_optD.sh) lands the previously-deferred Option D: worker admission_mode authoritative for direct_append + seed + reseed, bypassing replay's local _decode_session_soft_cap. Key findings now documented: - errors collapse from 9-10% to 0.2% (mooncake timeouts gone) - session-cap fallback rises 33-35% -> 46-51% — D's true KV pool is the binding constraint, not replay's estimator; v4's "low fallback" was hiding capacity overruns as transfer-timeout errors - direct-to-D subset latency unchanged from v4 (admission overhead negligible) - new bottleneck: D's physical KV pool — points v6 at prefill backup release timing, priority eviction tuning, chunked seed, cross-D session migration, and real RDMA Also adds a 5th lesson on errors-vs-fallback reciprocity and updates the code index with the v5 endpoint extension and new CLI knobs. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 16:13:25 +08:00
kzlin	6e5ed8da80	feat(kvc): Option D - delegate seed/reseed admission to D worker v4 (cap=16) saw 35% session-cap fallback because the local soft_cap min(16, usable / target) evaluates to 1-2 for large agentic inputs. The cap was hit not because D was full but because replay's heuristic underestimated capacity. This change makes worker admission_mode authoritative for ALL paths: SGLang side: - io_struct.py: DirectAppendAdmissionReqInput gains a `mode` field ("direct_append" \| "seed", default "direct_append" preserves prior behavior). - scheduler.py:admit_direct_append: when mode == "seed", skip the resident-on-D requirement and run the same capacity check + LRU eviction (maybe_trim_decode_session_cache) that direct_append uses. This lets D atomically decide if a new session can be admitted based on actual token_to_kv_pool_allocator state. Replay side (replay.py): - _query_decode_direct_admission gains a `mode` parameter. - _reserve_decode_session_capacity: in worker admission_mode, the seed/reseed branch now queries D with mode="seed" and trusts the result, instead of estimating capacity from the residency snapshot. - _should_admit_new_decode_session: in worker mode, skip the local soft_cap pre-check and let D decide. Same-D session fast-path is preserved. Effects: - Local hardcoded cap of 16 is bypassed under worker mode; D's real KV pool size is the only constraint. - LRU eviction runs in D's process atomically with admission, so starvation (the v3 bimodal "lucky vs starved sessions" pattern) should resolve. scripts/sweep_tp1_v5_optD.sh added to run the same 1P7D / 2P6D configs as v4 with the new admission path. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-28 23:40:03 +08:00
kzlin	74194e660a	docs: v4 final results, error analysis, and updated journey Add v4 sweep results and post-mortem analysis showing: - direct-to-D path: 54.3% (1P7D) / 58.0% (2P6D) of requests now use KVC cleanly. P50=0.5s and TTFT P50=0.043s; this path beats baseline 8DP across the board (P50 -24%, TTFT P50 -54%, TTFT P90 -79%). - Overall vs baseline (errors+truncated excluded): v4 2P6D P50=0.85s vs baseline 0.66s (28% slower). Reason is not errors -- 35% of requests still hit fallback-large-append-session-cap, where capacity-based cap = usable_tokens / target_tokens evaluates to 1-2 (not 16) for large agentic inputs. - 9-10% errors on KVC variants are mooncake TCP transfer timeouts, not SGLang logic bugs. Prefill log shows "Failed to send kv chunk ... 32s timeout ... session not alive". Errors concentrate in turn>=31 (large inputs) after run >44.8%. Track: - docs/KVC_DEBUG_JOURNEY_V1_TO_V4.md: append v4 results table, per-mode breakdown, and error root cause. - scripts/analysis/{analyze_v3,analyze_v4,analyze_errors,compare_no_error}.py - outputs/qwen3-30b-tp1-v{3,4}/exp_summary.json (force-added, small JSON; metrics.jsonl excluded due to size). - outputs/qwen3-30b-tp1-v{3,4}*/sweep_results.txt Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-28 23:34:01 +08:00
kzlin	c9d350b372	docs: KVC v1-v4 debug journey + raise session soft_cap to 16 Document the iterative debugging from v1 (broken KVC) through v4 (routing fixed + session cap raised), with code-level analysis of the two main bugs encountered: 1. v2 root cause (mis-diagnosed previously as `allow_local_prefill`): `--policy default` for KVC mechanism caused replay's round-robin policy and the PD router's round-robin to diverge, sending requests with `session_params` to a D worker that did not have the session open. Resulted in 56-61% truncation with finish_reason "session id X does not exist". Fix: use `--policy kv-aware` (sweep_tp1_v3_kvaware.sh) so replay emits `x-smg-target-worker` and PD router uses consistent_hashing. 2. v3 new bottleneck: `pd-router-fallback-large-append-session-cap` dominated 52-65% of requests. Root cause was hardcoded `min(4, ...)` in `_decode_session_soft_cap`. With 7 D workers x 4 sessions = 28 slots for 52 trace sessions, ~24 sessions starved permanently (bimodal direct-to-D rate of 0% or 99%). Fix: raise the cap to 16 (replay.py). Also includes the v3 finding that direct-to-d-session path P50=0.495s and TTFT P50=0.043s already beats the 8-way DP baseline (0.65s/0.093s) - the KVC core mechanism works when fallback paths are avoided. Files: - docs/KVC_DEBUG_JOURNEY_V1_TO_V4.md: full journey + code location index - docs/SWEBENCH_EXPERIMENT_{PROGRESS,RESULTS}.md: prior session notes - scripts/sweep_tp1_v{2,3,4}*.sh: experiment driver scripts - src/agentic_pd_hybrid/replay.py: cap 4 -> 16, audit fields - src/agentic_pd_hybrid/pd_router.py: strip session_params from prefill - src/agentic_pd_hybrid/metrics.py: truncated_request_count Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-28 21:10:41 +08:00
Gahow Wang	e9062b1d6e	Document PD baseline comparison	2026-04-25 17:29:27 +00:00
Gahow Wang	c928c7db23	Add transfer queue admission knobs	2026-04-25 17:29:15 +00:00
Gahow Wang	fe583fb413	Document kvcache-centric experiment progress	2026-04-25 16:01:31 +00:00
Gahow Wang	13bb31a446	Add kvcache-centric profiling and admission controls	2026-04-25 16:00:52 +00:00
Gahow Wang	08b13d22bc	docs: rewrite project docs in concise chinese	2026-04-24 12:41:52 +00:00
Gahow Wang	5bdc0ed4f0	docs: document sglang maintenance workflow	2026-04-24 12:31:32 +00:00
Gahow Wang	b8e6f13c20	feat(sglang): support decode session cache admission	2026-04-24 12:30:41 +00:00
Gahow Wang	bded08301f	chore: vendor sglang v0.5.10 snapshot	2026-04-24 12:29:36 +00:00
Gahow Wang	78f0d15221	docs: document project design and status	2026-04-24 12:17:55 +00:00
Gahow Wang	4bca741f32	feat: add agentic pd hybrid benchmark prototype	2026-04-24 12:17:46 +00:00