diff --git a/.gitignore b/.gitignore
index 9866984..94b93e7 100644
--- a/.gitignore
+++ b/.gitignore
@@ -13,6 +13,9 @@
 /third_party/llama.cpp/models/
 *.gguf
 
+# Claude Code runtime state
+/.claude/
+
 # Benchmark output + fetched datasets (transferred to GPU host, not committed)
 /bench-out/
 /tools/bench/data/
diff --git a/README.md b/README.md
index 9666e69..703e17a 100644
--- a/README.md
+++ b/README.md
@@ -144,16 +144,53 @@ HF_ENDPOINT=https://hf-mirror.com python3 -m tools.bench.fetch_datasets
 - `docs/00-roadmap.md`：总体路线图与各 Phase 设计
 - `docs/01..15-*.md`：CUDA FFI / Tensor / GEMM / Attention / KV cache / 性能优化等每个 Phase 的设计文档
 - `docs/16-llama-cpp-comparison.md`：llama.cpp 对比基准的设计
-- `docs/benchmarks/`：各阶段的 benchmark 报告
+- `docs/17-tensor-parallelism.md`：张量并行（TP）设计
+- `docs/18-pipeline-parallelism.md`：流水线并行（PP）设计
+- `docs/benchmarks/`：各阶段的 benchmark 报告（含 `pp-sweep.md`）
+
+## 多卡并行（TP / PP）
+
+单机多卡，复用 NCCL（crate `xserv-distributed`）。两种切法正交、二选一：
+
+- **张量并行 `--tp N`**：按 head / 中间维切每一层，层内用 AllReduce 聚合（每 token `2·层数` 次）。
+- **流水线并行 `--pp N`**：按层切成 N 段，相邻段间用 NCCL **P2P** 传 hidden state（每 token 仅 `N-1` 次），
+  通信量远小于 AllReduce，对无 NVLink 的 PCIe 更友好。
+
+```bash
+# 组内 GPU 0-3：4 卡张量并行 / 4 卡流水线并行
+CUDA_VISIBLE_DEVICES=0,1,2,3 ./target/release/xserv-server /path/to/qwen3-8b --tp 4
+CUDA_VISIBLE_DEVICES=0,1,2,3 ./target/release/xserv-server /path/to/qwen3-8b --pp 4
+```
+
+**PP 实测**（dash5，Qwen3-8B BF16，单流贪心；每卡显存为权重+最小 KV 池）：
+
+| 配置 | TTFT | TPOT | tok/s | 每卡显存 |
+|------|------|------|-------|----------|
+| 单卡 | 33ms | 17.4ms | 57.5 | 24.0 GB |
+| PP=2 | 36ms | 18.1ms | 55.3 | 11.6 / 13.6 GB |
+| PP=4 | 36ms | 17.9ms | 55.8 | 7.3 / 5.3 / 5.3 / 9.4 GB |
+
+**质量对比**（AIME 2025 30 题 + GSM8K 30 题，贪心，xserv 在 GPU 0-3、llama.cpp 在 GPU 4-7 并行）：
+
+| 引擎 | PP | AIME | GSM8K |
+|------|----|------|-------|
+| xserv | 1/2/4 | 8 / 7 / 7 (/30) | 29/30 (96.7%) 全部一致 |
+| llama | 1/2/4 | 7 / 7 / 7 (/30) | 29/30 (96.7%) 全部一致 |
+
+正确性：hidden state 跨段是 **bit-exact BF16 P2P 拷贝**，PP=4 输出与单卡逐字节一致（用「单卡×2 vs
+PP=4×2」对照确认——单卡自身因 cuBLAS 非确定性 run-to-run 会变，而 PP=4 可复现且落在某次单卡轨迹上）。
+GSM8K 12 个格子全是 29/30，xserv 与 llama.cpp 完全一致；AIME 的 ±1 是长生成下贪心对 GEMM 抖动的敏感，
+非 PP 或引擎效应。**收益在显存**（每卡权重+KV ≈ 1/N）；v1 为串行流水线，单流 TPOT 基本持平、不优于单卡，
+真正的吞吐提升需后续做 microbatch / 1F1B 重叠。完整数据见 `docs/benchmarks/pp-sweep.md`。
 
 ## 路线图（节选）
 
-已完成 Phase 0–15：CUDA 基础设施 → Tensor → GEMM → Transformer kernels → Attention →
+已完成 Phase 0–18：CUDA 基础设施 → Tensor → GEMM → Transformer kernels → Attention →
 模型加载 → 分词器 → GPT-2 → KV cache → Qwen3-8B → Paged Attention → 连续批处理 →
-HTTP API → Flash Attention 2 → 性能优化；并在此基础上加入了 **llama.cpp 对比基准**
-与 **KV CPU 换出** 等基础设施。
+HTTP API → Flash Attention 2 → 性能优化 → **张量并行（TP）** → **流水线并行（PP）**；
+并加入了 **llama.cpp 对比基准** 与 **KV CPU 换出** 等基础设施。
 
-后续方向：投机解码（speculative decoding）、张量并行（TP，多卡）、量化（FP8 / INT8）、多模态。
+后续方向：PP microbatch/1F1B 流水线重叠（吞吐收益）、2D TP×PP、投机解码、量化（FP8 / INT8）、多模态。
 
 ## 许可
 
diff --git a/docs/18-pipeline-parallelism.md b/docs/18-pipeline-parallelism.md
new file mode 100644
index 0000000..aeee592
--- /dev/null
+++ b/docs/18-pipeline-parallelism.md
@@ -0,0 +1,151 @@
+# Phase 18: Pipeline Parallelism (PP)
+
+> 目标：在单机多卡上做 **流水线并行**，把 Qwen3-8B 的 **层** 切成 `P` 段（stage），
+> 每张卡只持有连续的一段层（+ stage0 的 `embed_tokens`、最后一段的 `norm`/`lm_head`），
+> 激活（hidden state）在相邻 stage 之间用 **NCCL P2P send/recv** 传递。
+> 与 TP（按 head / 中间维切，每层 2 次 AllReduce）互补：PP 通信量小（每 token 仅 `P-1`
+> 次点对点传 `[tokens, hidden]`），KV 与权重按 **层** 降到约 1/P。
+> 先做 **PP=2 / 4（组内）**，正确性优先。
+
+## 1. 硬件约束（dash5）
+
+- 8× RTX 5090（32GB，SM120），**无 NVLink**，纯 PCIe Gen5。
+- 拓扑：GPU 0–3 一组、4–7 一组，组内 `PHB`（同 host bridge，可 P2P），跨组 `NODE`。
+- **PP 同样建议在组内**（0–3 或 4–7）：虽然 PP 的通信量远小于 TP，但 P2P 仍走 PCIe，
+  跨组延迟更高。PP=2/4 用 0–1 / 0–3。
+- 相比 TP：TP 每 token `2·layers = 72` 次 AllReduce（延迟主导）；PP 每 token 仅
+  `P-1` 次 send/recv，每次 `[tokens, hidden]` BF16（decode batch=1 时 8KB）。
+  **PP 对慢互联（PCIe / 无 NVLink）更友好**，这是在 dash5 上做 PP 的主要动机之一。
+
+## 2. 切分方案（layer-wise）
+
+Qwen3-8B：`hidden=4096`、`num_heads=32`、`num_kv_heads=8`、`head_dim=128`、
+`intermediate=12288`、`layers=36`、`vocab=151936`。`36` 能被 `2/4` 整除（PP=3/6 需处理余数，
+本阶段先要求 `layers % P == 0`）。
+
+设 stage 数 `P`，本 stage = `s`，每段 `L = layers / P` 层，本段持有全局层
+`[s·L, (s+1)·L)`：
+
+| 组件 | 持有者 | 说明 |
+|------|--------|------|
+| `embed_tokens` `[vocab, hidden]` | **仅 stage 0** | token → hidden |
+| transformer block `i` 的全部权重 | 持有 `i` 的那个 stage | 不切 head / 中间维（与 TP 正交） |
+| 该层 KV cache | 持有 `i` 的那个 stage | **每卡 KV 降到约 1/P** |
+| 最终 `norm` `[hidden]` | **仅最后一段** | |
+| `lm_head` `[vocab, hidden]` | **仅最后一段** | hidden → logits |
+
+- 注意力 / MLP 的层内计算 **完全不变**（不需要 AllReduce）：每个 stage 用它自己那几层
+  的完整权重、完整 head 做 forward。PP 与 TP 正交，可叠加（本阶段不实现 TP×PP）。
+- **RoPE** 用全局绝对 position，每个 stage 的 `RopeCache` 完全相同（按 position 索引），
+  各 stage 独立做，无需通信。
+- **每个 stage 一个独立的 `PagedKVCache`**，层数 = 本段层数 `L`（不是 36）。forward 时
+  按「本段内的局部层号 `0..L`」索引 cache —— 与单卡代码完全一致，只是 `self.layers`
+  只装了本段的层。实现技巧：给 cache 传一个 `num_hidden_layers` 改写成 `L` 的 config 克隆，
+  **无需改 `PagedKVCache`**。
+
+### 通信点
+- prefill：stage `s` 算完本段层，得到 `[S, hidden]` → **send 给 `s+1`**；`s+1` recv 后接着算。
+- decode：同理传 `[B, hidden]`（batch=1 时 `[1, hidden]`）。
+- 每 token 共 `P-1` 次 send/recv；最后一段算出 logits 并采样。
+- 采样得到的 token id（一个 `u32`）由 **最后一段经线程内 channel 回传给 stage0**
+  （同进程多线程，无需走 NCCL）。
+
+## 3. 进程 / 线程模型
+
+沿用 TP 的 **单进程、多线程**：每个 stage 一个 OS 线程，线程启动时 `cudaSetDevice(stage)`。
+- **stage 0 = 协调者（coordinator）**，跑在调用线程上：持有 scheduler、tokenizer、HTTP
+  response sender、停止判定（eos / max_tokens）与「下一步输入 token」。
+- **stage 1..P-1 = worker 线程**：从控制 channel 收命令（Register/Prefill/Decode/Free/Shutdown），
+  每步 `recv` 上游 hidden → 跑本段层 → `send` 给下游；最后一段 `head`+采样 → 把 token 回传 stage0。
+- 控制信息（命令、采样参数、token id）走 `mpsc`（极小）；**重活（hidden 张量）走 NCCL P2P（GPU↔GPU）**。
+
+> **v1 串行语义**：一次处理一个请求、一次一个 token，流水线每步「灌满又排空」
+> （stage0 decode 第 `t+1` 步依赖最后一段第 `t` 步采出的 token）。这保证 **正确性**，
+> 并拿到 TTFT/TPOT 与每卡显存；**throughput 的真正收益来自 microbatch/请求级流水线
+> 重叠（1F1B）**，列为后续工作（见 §7）。
+
+执行流（每请求）：
+```
+coordinator                     worker s (1..P-1)            last stage (P-1)
+─────────────                   ─────────────────            ────────────────
+broadcast Register(slot)        cache.register(slot)         cache.register(slot)
+broadcast Prefill{n,slot,samp}
+ x=embed(prompt)
+ x=layers_prefill(x,slot)
+ send x → stage1                recv x ← s-1
+                                x=layers_prefill(x,slot)
+                                send x → s+1 ───────────────► recv x ← P-2
+                                                             x=layers_prefill(x,slot)
+                                                             logits=head(x); next=sample
+ next ◄────────────── token channel ◄──────────────────────  token_tx.send(next)
+ stream(next); loop Decode{slot} 直到 eos/length
+broadcast Free(slot)            cache.free(slot)             cache.free(slot)
+```
+
+## 4. 通信库：NCCL P2P
+
+复用 `xserv-distributed`（已有 NCCL FFI + `TpContext`/AllReduce），新增：
+- FFI：`ncclSend(sendbuff, count, dtype, peer, comm, stream)`、
+  `ncclRecv(recvbuff, count, dtype, peer, comm, stream)`。
+- `PpContext`：与 `TpContext` 同样的 `ncclCommInitRank`（一个 comm 跨 `P` 个 stage），
+  外加 `send_bf16_ptr(ptr, count, peer)` / `recv_bf16_ptr(ptr, count, peer)`，在 **null
+  stream** 上发起（与模型 kernel 同流，天然有序）。
+- 线性流水线无死锁：stage0 只 send、最后一段只 recv、中间段「先 recv 上游、再 send 下游」，
+  依赖链无环，从头解锁。每个 stage 在 send/recv + 本段计算后 `synchronize()`，
+  确保 NCCL 读完发送缓冲再复用/释放（v1 串行下成本可接受）。
+
+> **决策点**：和 TP 一样，collective/P2P 先用 NCCL 把 PP 跑通拿正确性与基线；
+> 手写 P2P（PCIe 上的 cudaMemcpyPeer）作为后续学习项。
+
+## 5. 权重分片加载
+
+`Qwen3::from_weights_pp(config, weights, stage, num_stages, device)`：
+- 只把全局层 `[s·L, (s+1)·L)` 搬到本 stage 的 GPU（其余层的权重直接 drop，不占显存）。
+- `embed_tokens`：仅 stage 0 加载；其余 stage 放一个 1×1 占位张量（forward 用 `is_first_stage`
+  守卫，永不触碰）。
+- `norm`/`lm_head`：仅最后一段加载；其余放占位。
+- head 不切（不做 TP），所以 `local_num_heads = num_heads`、`local_num_kv_heads = num_kv_heads`。
+
+每卡显存 ≈ `权重(transformer 1/P) + KV(1/P) + (stage0: embed) + (last: norm+lm_head)`。
+对 Qwen3-8B：transformer 层约 14GB，PP=2 每卡约 7GB 层权重 + embed 或 lm_head（各 ~1.2GB）。
+
+## 6. 实现步骤（逐步可验证）
+
+1. **P18.1 — `xserv-distributed` P2P**：`ncclSend/Recv` FFI + `PpContext`。
+   验收：2 卡，rank0 send 已知向量、rank1 recv，校验一致（`tests/sendrecv.rs`）。
+2. **P18.2 — 分段权重加载**：`from_weights_pp`，每 stage 只持有本段层 + 该有的 embed/head。
+   验收：各 stage 层数 = `L`、显存约 1/P（+ embed/head）。
+3. **P18.3 — stage forward**：`embed` / `forward_layers_prefill` / `forward_layers_decode` /
+   `head`，每段独立 KV cache。
+   验收：**PP=1 与单卡 `forward_*_paged` 逐 token 一致**（同一条代码路径退化）。
+4. **P18.4 — PP engine + `--pp N`**：多线程 stage workers + NCCL 传递 + stage0 协调。
+   验收：`--pp 2/4` 端到端可服务；**greedy 输出与单卡（PP=1）逐 token 一致**；
+   用现有 llama.cpp bench 跑正确性（GSM8K/AIME）；测 PP=1/2/4 的 TTFT/TPOT/每卡显存。
+
+## 7. 预期与风险
+
+- **显存**：每卡 transformer 权重 + KV ≈ 1/P，这是 PP 的主要收益（可上更大模型 / 更长 context）。
+- **单流吞吐**：v1 串行无 stage 重叠 → 单流 tok/s **不会超过单卡**（多一份 P2P + sync 开销，
+  可能略低）。这是 PP 的本质：**没有 microbatch 重叠就没有加速**。诚实记录实测，并与
+  llama.cpp 的 `--split-mode layer`（同样是层切流水线、单序列也串行跨卡）对比 —— 两者单流
+  都应≈单卡。
+- **真正的 throughput 收益**（后续）：请求级 / microbatch 流水线（1F1B），让 stage 间重叠：
+  stage1 算 microbatch A 时 stage0 算 B。需要把 scheduler 改成跨 stage 连续批处理。
+- **风险**：NCCL 多线程 init 同步；send 缓冲生命周期（必须 sync 后再复用）；
+  `layers % P != 0` 的余数分配（本阶段先约束整除）；与 CUDA Graph decode 的结合（先走非 graph 路径）。
+- 正确性优先：先 PP=1 等价（逐 token 对齐），PP=2/4 与单卡对齐，再谈性能。
+
+## 8. 与 llama.cpp 的对比口径
+
+- **xserv**：`--pp N`，`CUDA_VISIBLE_DEVICES=0..N-1`。
+- **llama.cpp**：`-sm layer`（默认即层切流水线）+ `--tensor-split` 均分层，`CUDA_VISIBLE_DEVICES=0..N-1`。
+  （对照 TP 用的是 `-sm row`。）
+- 指标：正确性（GSM8K / AIME exact-match）、单流 TTFT/TPOT、并发吞吐、每卡 VRAM。
+- 复用 `tools/bench/runner.py` 与 `run_pp_parallel.sh`（仿 `run_tp_parallel.sh`）。
+
+## 9. 不在本阶段范围
+
+- TP×PP 混合（2D 并行）、跨组 / 多节点。
+- microbatch / 1F1B 流水线重叠（throughput 收益，后续）。
+- vocab-parallel embedding / lm_head。
+- `layers % P != 0` 的非均匀切分；与 CUDA Graph decode 结合。
diff --git a/docs/benchmarks/pp-sweep.md b/docs/benchmarks/pp-sweep.md
new file mode 100644
index 0000000..5f1f5cd
--- /dev/null
+++ b/docs/benchmarks/pp-sweep.md
@@ -0,0 +1,118 @@
+# PP sweep — xserv vs llama.cpp (Qwen3-8B BF16, 8×RTX 5090)
+
+Pipeline parallelism (layer split), verified end-to-end on dash5. Qwen3-8B BF16,
+greedy, single stream, no NVLink (hand-off / split traffic over PCIe Gen5).
+xserv `--pp N` puts stage `s` on GPU `s` and hands the hidden state stage→stage
+over NCCL P2P; llama.cpp uses `-sm layer` (its default pipeline split) over N GPUs.
+
+## Single-stream latency + per-GPU VRAM (measured, `--max-seq-len 2048`)
+
+Measured strictly sequentially, one server at a time, each config gated on a real
+successful generation (so VRAM snapshots are post-load). Driver:
+`tools/pp_final.sh`.
+
+| engine | PP | TTFT_ms | TPOT_ms | tok/s | per-GPU VRAM (MiB) |
+|--------|----|---------|---------|-------|--------------------|
+| xserv  | 1  | 33.2    | 17.39   | 57.5  | 24010 |
+| xserv  | 2  | 35.9    | 18.07   | 55.3  | 11580, 13632 |
+| xserv  | 4  | 36.1    | 17.91   | 55.8  | 7298, 5250, 5250, 9350 |
+| llama  | 1  | 133.3   | 9.38    | 106.7 | 15604 |
+| llama  | 2  | 131.4   | 9.10    | 109.9 | 7862, 8494 |
+| llama  | 4  | 161.2   | 8.88    | 112.6 | 4476, 4090, 4090, 5108 |
+
+(xserv VRAM with `XSERV_MAX_KV_BLOCKS=160` so the number is weights + a minimal
+KV pool. `tok/s = 1000 / TPOT`. This latency probe's TTFT differs from the
+quality-suite TTFT below because the suite includes scheduler/HTTP overhead.)
+
+## Correctness — PP is numerically exact
+
+The hidden-state hand-off between stages is a bit-exact BF16 P2P copy and each
+stage runs the same kernels over its layers, so PP must reproduce the single-GPU
+result. Verified by byte-comparing generated text (greedy, temp 0), running each
+config **twice** to separate PP effects from run-to-run GEMM noise:
+
+| comparison | result |
+|------------|--------|
+| single run A == single run B | **DIFFER** (cuBLAS GEMM is not bit-reproducible run-to-run) |
+| pp4 run A == pp4 run B | **IDENTICAL** |
+| single run A == pp4 run A | **IDENTICAL** |
+| single == pp2 (single run each) | **IDENTICAL** |
+
+Takeaway: **single-GPU itself is non-deterministic** under greedy (a 1-ULP logit
+difference flips a late argmax and the suffix changes), so a one-shot single-vs-PP
+byte compare can spuriously "DIFFER". The 2×2 control shows PP=4 is *more*
+reproducible than re-running single-GPU, and it lands exactly on a single-GPU
+trajectory. NCCL P2P (`tests/sendrecv.rs`) and AllReduce (`tests/allreduce.rs`)
+unit tests pass.
+
+## Quality matrix — AIME 2025 (30) + GSM8K (30), greedy, both engines × PP=1/2/4
+
+Full measured matrix (`tools/bench/summarize_fullq.py`; raw in
+`bench-out/FULLQ_SUMMARY.txt`). Qwen3-8B BF16, thinking OFF, `max_seq_len 4096`.
+xserv on GPUs 0-3, llama.cpp on GPUs 4-7 (disjoint groups, run in parallel).
+
+| engine | PP | AIME 2025 | GSM8K | AIME mean_tok | TTFT_ms | TPOT_ms |
+|--------|----|-----------|-------|---------------|---------|---------|
+| xserv  | 1  | 8/30 (26.7%) | 29/30 (96.7%) | 2383 | 485 | 22.42 |
+| xserv  | 2  | 7/30 (23.3%) | 29/30 (96.7%) | 2367 | 457 | 22.55 |
+| xserv  | 4  | 7/30 (23.3%) | 29/30 (96.7%) | 2652 | 494 | 23.31 |
+| llama  | 1  | 7/30 (23.3%) | 29/30 (96.7%) | 2651 | 119 | 10.37 |
+| llama  | 2  | 7/30 (23.3%) | 29/30 (96.7%) | 2651 | 118 | 10.41 |
+| llama  | 4  | 7/30 (23.3%) | 29/30 (96.7%) | 2651 | 119 | 10.39 |
+
+Reading the matrix:
+
+- **GSM8K = 29/30 (96.7%) in every cell** — identical across both engines and all
+  PP levels. xserv's accuracy matches llama.cpp exactly on the same weights.
+- **AIME = 7/30 (23.3%) everywhere except xserv PP=1 (8/30)**. That single +1 is
+  the run-to-run greedy nondeterminism documented above (an AIME solution is
+  ~2400 tokens; one late argmax flip changes one problem's outcome) — not a PP or
+  engine effect. AIME accuracy is low because this is an 8B model with thinking
+  disabled; the point here is the *cross-engine / cross-PP agreement*, which holds.
+- **TPOT is flat across PP** for both engines (xserv 22.4→23.3 ms, llama
+  10.3→10.4 ms), reconfirming PP doesn't slow single-stream decode. The ~2.2×
+  TPOT gap to llama.cpp is the single-GPU gap (`llama-cpp-comparison.md`),
+  orthogonal to PP.
+
+## Takeaways
+
+- **Memory is the win.** Per-GPU weights+KV scale ~1/P: xserv 24.0 GB (1 GPU) →
+  ~11–14 GB (PP=2) → ~5–9 GB (PP=4); llama 15.6 → ~8 → ~4–5 GB. The two end
+  stages sit higher (stage 0 holds `embed_tokens`, the last stage `norm`+`lm_head`,
+  ~1.1 GB each). This is what PP buys: a model / context that does not fit on one
+  card fits across P.
+- **Single-stream latency is flat, not faster.** v1 PP is serial across stages
+  (no microbatch overlap): per-token latency = sum of all stages' compute +
+  (P-1) P2P hops + a blocking sync per stage. The `[1, hidden]` BF16 hop (8 KB)
+  over PCIe is cheap relative to per-token compute, so TPOT is ~constant across P.
+  PP does **not** speed up single-stream decode; it trades (almost no) latency for
+  large memory headroom.
+- **Quality is preserved and matches llama.cpp.** GSM8K 96.7% in all 12 cells;
+  AIME within the greedy noise band. PP=1/2/4 agree, and xserv tracks llama.cpp.
+
+## Reproduce
+
+```bash
+./tools/sync-and-build.sh build
+# latency + VRAM + byte-exact correctness (writes bench-out/PP_FINAL.md):
+ssh <host> 'cd <repo> && bash tools/pp_final.sh'
+# determinism control (single×2 vs pp4×2):
+ssh <host> 'cd <repo> && bash tools/pp_diag.sh'
+# NCCL P2P + AllReduce unit tests:
+ssh <host> 'cd <repo> && cargo test -p xserv-distributed --release'
+# full quality matrix AIME-30 + GSM8K-30 (xserv 0-3 serial; or parallel w/ llama 4-7):
+ssh <host> 'cd <repo> && bash tools/pp_quality_full.sh'   # xserv+llama serial, GPU 0-3
+ssh <host> 'cd <repo> && bash tools/pp_llama_47.sh'        # llama on GPU 4-7 (parallel)
+python3 tools/bench/summarize_fullq.py bench-out
+```
+
+## Next (where PP actually raises throughput)
+
+- **Microbatch / 1F1B overlap**: while stage 1 runs microbatch A, stage 0 runs B.
+  This is the only thing that turns PP into a *throughput* win; v1 is serial, so
+  P GPUs give 1 GPU's single-stream rate (but P× the memory headroom / batch room).
+- Persistent per-stage recv buffers (drop the per-token CPU alloc + H2D) and
+  event-based ordering instead of a full device sync per hop.
+- 2D TP×PP, and `layers % P != 0` non-uniform splits.
+
+🤖 Generated with [Claude Code](https://claude.com/claude-code)