PXN Send-Side Engine 分阶段开发计划

1. 目标与开发契约

目标文件：/Users/joker/Desktop/project/baai/sglang/liuda/5-pxn-send-side-engine-dev.md。进入执行模式后的第一步，将本计划写入该文件。

实现基线：

Mooncake：079267fc（Commit 3）
SGLang：5899674504
仅修改 Mooncake/mooncake-transfer-engine
SGLang、decode、Python/C++ 公共传输 API 零改动
不修改 RDMA WorkerPool 的 post/poll 分工
MC_PXN_ENABLE=0 时不创建 Pump、HBM arena、shm 或 CUDA IPC 映射
先提交 RFC，再开始 Mooncake 代码开发
一个 PR、多个顺序 commit
Codex 每阶段只实现和测试，不执行 git add、git commit、git push，由用户手工提交

RFC 建议标题：

[RFC][TransferEngine] Send-side PXN relay for rail-aligned RDMA

以下现有无关文件不得带入提交：

Mooncake 的 mooncake-transfer-engine/claude.md
SGLang 中旧 Python relay/ALLOC-COPY 版本的 test_mooncake_pxn.py
.claude/、response.md 及其他用户改动

2. 冻结后的实现设计

2.1 线程与调用关系

上层调用线程只负责分类、复制请求元数据、建立本地完成对象和投递 MPSC，不扫描门铃，不等待 lane credit。

SGLang worker / TE 调用线程
  └─ MultiTransport::submitTransfer()
       └─ PxnRdmaTransport::submitTransferTask()
            ├─ 分类并复制 TransferRequest
            ├─ 同轨、READ、CPU、rail 不唯一
            │    └─ RdmaTransport::submitTransferTask() 直接发送
            └─ 所有可走 PXN 的跨轨 CUDA WRITE，不按大小豁免
                 ├─ PieceBuilder：按 8MiB / 4096 plan 双上限切片
                 ├─ 为全部 Piece 建立本地 Slice 和 task 计数
                 ├─ cudaEventRecord(cudaStreamPerThread)
                 └─ 加入 per-lane mutex + deque MPSC，随后返回

PxnPump 是唯一新增线程，同时推进 sender 和 relay。

PxnPump
  ├─ reapOutbound()
  │    └─ 读取远端 Completion
  │         ├─ markSuccess/markFailed 本地 logical Slice
  │         └─ 推进 reaped_seq，之后才允许复用 slot
  ├─ reapInbound()
  │    └─ 查询 relay child RDMA batch
  │         ├─ 写 Completion{seq,status}
  │         ├─ 连续推进 completed_seq
  │         └─ freeBatchID()
  ├─ progressInbound()，按 7 条 lane round-robin
  │    └─ 门铃 → 校验 Desc → session name 解析本地 SegmentID
  │         └─ 建立由 InflightRecord 持有的 child tasks
  │              └─ pxn_transport->RdmaTransport::submitTransferTask()
  │                   └─ Pump 调用栈上切成 64KiB RDMA Slice
  │                        └─ 既有 WorkerPool
  │                             ├─ thread 1..N → ibv_post_send
  │                             └─ thread 0 → ibv_poll_cq
  ├─ progressOutbound()，按 7 条 lane round-robin
  │    └─ ready event → credit → desc → CUDA copy → doorbell
  └─ 无进展时 condition_variable::wait_for(50µs)

Pump 每轮顺序固定为：

限频发布 heartbeat。
收割 outbound completion。
收割 inbound RDMA child batch。
round-robin 推进入方向，每 lane 每轮至多一票。
round-robin 推进出方向，每 lane 每轮至多一票。
推进尚未发布的直发 fallback。
无工作时休眠。

credit 不足只能跳过当前 lane，不能阻塞 Pump，否则双向 rank 可能互相等待。

Relay 不调用继承的 RdmaTransport::submitTransfer()，因为它内部会虚调用 submitTransferTask()，可能重新进入 PXN override。Relay 自行建立 child BatchDesc/TransferTask，然后限定调用：

pxn_transport->RdmaTransport::submitTransferTask(child_tasks);

2.2 Transport 接入方式

采用：

class PxnRdmaTransport : public RdmaTransport

约束：

继承现有 install、MR 注册、batch、status 和连接管理。
直接路径限定调用 RdmaTransport::submitTransferTask()。
Relay child batch同样限定调用基类 task 方法。
getTransferStatus() 继续使用现有 task/slice 计数，不读取 shm。
不给 Transport::Slice::union 增加 PXN 字段。
MultiTransport 仅在 Linux/CUDA/RDMA 编译条件满足且 MC_PXN_ENABLE=1 时构造子类。
协议名仍为 rdma。
PXN 初始化失败后保留普通 RDMA 能力并安全直发。

2.3 Arena 与 lane 布局

v1 固定 7 条物理 lane，只对应另外 7 个 rank，不分配 self lane。允许调整的布局参数只有编译期的 K 和 S：

K：每条 lane 的 slot 数，当前为 18
S：每个 slot 的大小，当前为 8MiB
StagingArena 大小只能由 LANE_COUNT × K × S 推导，不提供独立容量配置或环境变量

精确布局：

LANE_COUNT = 7
K = 18
S = 8MiB
 
每 lane 数据区：
18 × 8MiB = 144MiB
 
7 lane 数据区：
7 × 144MiB = 1008MiB
 
StagingArena：
1008MiB

arena 不提供独立容量配置。调整编译期的 slot 数或 slot 大小时，kRequiredArenaSize = LANE_COUNT × K × S 随之变化；运行时始终按该值分配。

地址计算固定为：

lane_base = arena_base + lane_id × K × S
slot_base = lane_base + slot_id × S
slot_id   = (seq - 1) % K

当前 K=18、S=8MiB，因此 lane stride 为 144MiB。禁止因为当前值误用位运算取模。

少于 8 个本机 rank 时，未分配的 lane 保持空闲；发现超过 7 个合法 peer 时，该进程禁用 PXN 并直发，不做不确定的部分映射。

2.4 Piece 与任意长度流水

当前参数：

Piece 最大为 S，当前 8MiB
每条 lane 有 K 个 slot，当前 18
MAX_PLAN_COUNT = 4096
跨轨路由不设大小阈值
session name 最大 256 字节

以一次 5GiB 连续传输为切片边界测试示例：

5GiB TransferRequest
  └─ PieceBuilder → 640 × 8MiB Piece
       └─ 同一 sender→relay lane 最多 18 个已发布 Piece 在飞
            ├─ seq 1..18 使用 slot 0..17
            ├─ sender 收割 seq 1 到本地 Slice
            ├─ reaped_seq 前进到 1
            └─ seq 19 才允许复用 slot 0

cudaMemcpyBatchAsync 一次只处理一个不超过 S 的 Piece，所有目标指针都位于该 Piece 对应的单个 staging slot。当前 S=8MiB；不会一次搬运整个 5GiB，也不要求 arena 同时容纳完整请求。

Piece 同时受两个上限约束：

piece_len <= S
plan_count <= 4096

任一上限到达立即切 Piece。单个 block 大于 S 时按地址范围拆分。

一个 Piece 可以包含多个原始 TransferTask 的数据段。每个数据段建立一个 logical Slice；Piece 完成时完成其全部 Slice。所有 task 的 slice_count 必须在任何 Piece 入队前建立完成，防止 event-driven completion 提前结束 task。

2.5 CUDA 顺序

提交线程捕获 source-ready event，Pump 的 per-lane CUDA stream 执行：

cudaStreamWaitEvent(lane_stream, ready_event)
  └─ cudaMemcpyBatchAsync / CUDA <12.8 cudaMemcpyAsync fallback
       └─ cuStreamWriteValue64(doorbell, seq)

要求：

CUDA 12.8、CUDA 13 分别适配已有的两种 cudaMemcpyBatchAsync 签名。
CUDA <12.8 使用逐项 cudaMemcpyAsync。
使用 cudaMemcpySrcAccessOrderStream。
健康路径禁止 cudaStreamSynchronize 或逐 Piece CPU blocking。
不重构现有 NVLink transports；PXN 内实现小型内部 helper，沿用相同编译分支。
不支持 stream memory op 时，在 lane stream 记录 event，Pump 在 cudaEventQuery 成功后由 CPU release-store doorbell。
CUDA copy 或 doorbell enqueue 失败且门铃尚未提交时，quarantine lane并允许直发。
cuStreamWriteValue64 提交成功即视为已经发布，之后不得直发重放。

2.6 Ring ABI 与完成所有权

序列规则：

next_seq 初始为 1。
reaped_seq、published_seq、completed_seq 初始为 0。
slot 为 (seq - 1) % 18。
分配 seq 前要求 next_seq - reaped_seq <= 18。
K=18 不能安全跨越 uint64 回绕；seq 为 0 非法，达到 UINT64_MAX 后先排空 lane，再更换 epoch 并从 1 重新建立 ring。

共享结构采用固定宽度、cache-line 对齐 POD，并以 static_assert 固定 ABI：

Registry header：magic、version、header size、state、uid、pid、进程 start time、boot epoch、rail、arena 信息、CUDA IPC handle。
Desc：commit seq、piece length、plan count、session length、256 字节 session、最多 4096 个 {final_dst, len}。
Completion：commit seq、int32_t Mooncake 状态码。
Relay 校验 epoch、seq、长度、plan count、地址溢出和 sum(plan.len) == piece_len。
非法 Desc 写失败 Completion 并归还该票，不能永久泄漏 credit。
Desc payload 先写，commit seq 最后 release-store。
Completion status 先写，completion seq 最后 release-store。
Relay 可以保存乱序终态，但 completed_seq 只能连续推进。
Sender 先将结果落入本地 Slice，再推进 reaped_seq。
slot 复用只依赖 sender 已完成本地收割的 reaped_seq。
上层 getTransferStatus() 永远只读取本地 task/slice 计数。

2.7 Registry、rail 与生命周期

启用 PXN 必须显式提供：

MC_PXN_GROUP_ID=<1..64 字节的 [A-Za-z0-9_.-]+，且不为 . 或 ..>

缺失或非法时只警告一次并直发。

MC_PXN_GROUP_ID 是本机 /dev/shm Registry 的作业隔离命名空间。同一作业的本机 rank 必须配置相同值；它不是 rank ID、TP/PP group、decode session 或 rail 名。

Registry 位于 /dev/shm：

目录包含 uid 和 group id。
目录权限 0700，文件权限 0600。
记录 pid、进程 start time、随机 boot epoch、rail、IPC handle、lane assignment、heartbeat，以及 lane count、K、S 和派生 arena size。
完全初始化后才发布 READY。
flock 仅用于注册、lane 分配和 stale cleanup。
热路径不使用文件锁。
self identity 不参与 lane 分配。
每个 relay 最多为另外 7 个 rank 分配 7 条 inbound lane。
relay 重启产生新 epoch，旧 lane 和旧 IPC mapping 全部失效。
Registry 中的 lane count、K、S 或 ABI version 不一致时拒绝映射并直发。

rail 解析：

本进程必须只有一个可用本地 HCA。
目标 SegmentDesc 必须解析出唯一目标 HCA。
MC_PXN_RAIL_MAP 为 raw HCA 到 canonical rail 的 JSON object，不保存 rank 到 rank 的映射。
Registry 必须找到唯一、健康、非 self 且 canonical rail 匹配的 relay。
relay 侧同一目标 session 也必须解析出唯一 HCA。
任一步不唯一、缺失或不支持时直发。

初始化顺序：

完成普通 RdmaTransport::install()。
校验配置、CUDA IPC 和本机拓扑。
cudaMalloc(kRequiredArenaSize)。
通过基类注册本地 RDMA MR，update_metadata=false。
创建 shm control mapping；sender 侧 CUDA host mapping 在 Commit 3 建立。
导出 CUDA IPC handle。
发布 Registry READY。
启动 Pump。

关闭顺序：

停止接受新 PXN Piece。
唤醒 Pump。
在有界时间内收割或终结队列和 inflight。
join Pump。
销毁 stream/event，关闭 peer IPC。
注销 sender host mapping 和 shm。
取消 Registry 发布并 unlink。
注销 staging MR。
cudaFree staging arena。
进入基类析构。

2.8 配置接口

不增加公共方法，仅新增环境变量：

配置	默认值	规则
`MC_PXN_ENABLE`	`0`	关闭时零额外资源
`MC_PXN_GROUP_ID`	无	开启 PXN 时必须显式设置
`MC_PXN_INFLIGHT_DEPTH`	`4`	有效范围 1..64
`MC_PXN_RAIL_MAP`	空	JSON raw HCA→canonical rail
`MC_PXN_CREDIT_TIMEOUT_MS`	`1000`	只影响未发布 Piece
`MC_PXN_HEARTBEAT_TIMEOUT_MS`	`3000`	超时后停止发布新票

MC_PXN_INFLIGHT_DEPTH 是单个 relay PxnPump 全局共享的第二跳 RDMA Piece 上限，不是每 lane 的窗口，也不改变每 lane 的 18 个 slot。默认 4 表示同时最多有 4 个 Piece 已提交给第二跳 RDMA。

MC_PXN_RAIL_MAP 只做 HCA 名称标准化。实际 sender→relay rank 关系由 Registry 根据 canonical rail 动态选择。

v1 固定值：

heartbeat interval：100ms
Pump idle wait：50µs
lane count：7
slots/lane：18，可通过修改编译期 K 调整
slot size：8MiB，可通过修改编译期 S 调整
跨轨 CUDA WRITE：不设大小阈值
max plan count：4096
max session：256 字节

2.9 失败提交点

以 GPU doorbell 是否成功提交为所有权边界：

发布前：peer 缺失、credit timeout、ready/copy/doorbell enqueue 失败均可通过基类 RDMA 直发。
发布后：不得透明重放。
heartbeat 暂停只阻止新票，旧票等待明确 completion、epoch 失效或已有传输超时。
relay epoch 消失时，旧 epoch 已发布票失败；重新发现后只有新票使用新 epoch。
第二跳 RDMA 失败直接使原 logical Slice 失败。
fallback child batch 和 relay child batch必须持有自己的 TransferRequest vector 直到 freeBatchID()。
所有 fallback 只提交一次，并把结果映射回原 logical Slice。

3. 分阶段 Commit

Phase 0：RFC 与计划冻结

无 Mooncake 代码改动。

将本计划写入 liuda/5-pxn-send-side-engine-dev.md。
建立 Mooncake RFC。
RFC 明确 7 lanes、18 slots/lane、8MiB/slot，以及派生的 1008MiB arena。
计划文件不混入 Mooncake PR。

Commit 1：配置、协议和纯逻辑切片

[TransferEngine] Add PXN configuration and protocol primitives

实现：

配置解析、rail map、序列工具、Desc/Completion ABI。
固定 L=7、K=18、S=8MiB。
PieceBuilder 按 8MiB/4096 entry 双约束切分。
尚不接入运行路径。

测试：

kRequiredArenaSize = 7×18×8MiB = 1008MiB。
(seq-1)%18，明确验证不能按位取模。
S-1/S/S+1 切片边界。
5GiB 生成 640 个 Piece。
plan cap、session 超长、溢出、回绕前换 epoch、malformed Desc。

Commit 2：Registry、arena 与 7 条 staging lane

[TransferEngine] Add PXN registry and staging lanes

实现：

/dev/shm/mooncake-pxn-<euid>-<group_id>/ Registry、READY、epoch、PID start time 和 stale cleanup。
READY 通过同目录 renameat2(RENAME_NOREPLACE) 原子发布，.lock 保留且不删除。
7 条非 self inbound lane；lane owner 同时记录 sender process identity 和 sender epoch。
每 lane 18 个8MiB slot。
1008MiB HBM arena、local-only RDMA MR、CUDA IPC export/import。
sender 的 IPC mapping 与 lane lease 绑定；relay 撤销 READY 和 sender attach 使用同一 entry lock，lease 未清零时不释放 HBM。
初始化回滚和逆序析构。
尚不接入公开 RDMA 路径。

Commit 2 的实际初始化调用为：

手工构造 PxnRdmaTransport / 后续 Commit 5 的 rdma factory
  └─ PxnRdmaTransport::install()
       ├─ RdmaTransport::install()
       └─ LocalResources::Create()
            ├─ RailResolver::resolveUnique(active_hcas)
            ├─ Registry::Open(group_id)
            ├─ StagingArena::initialize()
            │    ├─ cudaMalloc(1008MiB)
            │    ├─ RdmaTransport::registerLocalMemory(false, false)
            │    └─ cudaIpcGetMemHandle()
            └─ RegistryRegistration::publish() ← READY 最后发布

peer staging 映射调用为：

LocalResources::mapPeer(entry)
  ├─ Registry::mapPeer() → mmap relay control block
  ├─ PeerMapping::claimLane(sender identity, sender epoch)
  ├─ PeerMapping::attachArena() → 在 entry lock 内建立 lane lease
  └─ cudaIpcOpenMemHandle() → 取得 relay arena 在本进程的 VA

如果 RDMA MR 注册失败，现有基类接口无法证明所有 partial MR 均已注销，因此该 arena 必须 quarantine，不能冒险 cudaFree；进程退出时再由 CUDA/RDMA runtime 回收。

本阶段只保留 pxn_lifecycle_test 的两个 smoke case：完整 publish/map/lease/shutdown，以及初始化失败回滚。完整故障矩阵放到后续集成验收，不在 Commit 2 堆叠单测。

Commit 2 只编译资源 adapter，MultiTransport 仍构造普通 RdmaTransport；正常公开路径到 Commit 5 才切换为 PxnRdmaTransport。

Commit 3：Sender staging pipeline

[TransferEngine] Add the PXN sender pipeline

实现：

SenderPipeline 管理 peer lane，SenderLane 使用 mutex + deque 接收多生产者提交。
提交线程记录 cudaEventRecord(cudaStreamPerThread)，不等待 credit 或 CUDA copy。
Pump 侧接口为 reapOutbound() 和 progressOutbound()；本阶段不创建线程。
每 lane 独立 CUDA stream，按 ready fence → batch copy → doorbell 顺序发布。
CUDA 12.8、13 使用对应的 cudaMemcpyBatchAsync 签名，旧版本逐项调用 cudaMemcpyAsync。
支持 cuStreamWriteValue64 时由 GPU 写门铃；否则 event query 成功后由 Pump CPU release-store 门铃。
sender 收割 Completion 后先完成本地 logical Slice，再推进 reaped_seq 并归还 slot credit。
copy 或 doorbell 提交失败时 quarantine lane；门铃未发布的 Piece 使用限定基类 RDMA child batch fallback。
fallback 提交前先 query 原 ready event，未 ready 时留在 Pump 队列，禁止 RDMA 抢跑读取 source。
fallback ticket 持有自己的 TransferRequest vector，完成后再 freeBatchID()。
sender fallback已合并到 pxn_sender.cpp；它只把已经进入异步 sender queue、但尚未发布 doorbell的 Piece包装成自持有的 child batch，并限定调用现有 RdmaTransport::submitTransferTask()。
PxnRdmaTransport::install() 创建 sender pipeline，但不覆盖 submit、不启动 Pump，也不接入 factory。

Commit 3 的调用关系为：

未来 Commit 5 的 PxnRdmaTransport::submitTransferTask()
  └─ SenderLane::enqueue()
       ├─ recordReady(cudaStreamPerThread)
       └─ per-lane mutex + deque MPSC
 
未来 Commit 4 的 PxnPump
  ├─ SenderPipeline::reapOutbound()
  │    └─ Completion → logical Slice success/failure → reaped_seq
  └─ SenderPipeline::progressOutbound()
       ├─ credit 检查与 Descriptor commit
       ├─ cudaStreamWaitEvent(lane stream, ready event)
       ├─ cudaMemcpyBatchAsync / cudaMemcpyAsync
       └─ cuStreamWriteValue64 或 event query → CPU doorbell

测试：

pxn_sender_test 只保留两个 smoke case：MPSC + 18-slot credit/reap/reuse，以及 copy 完成后门铃与发布前 fallback。

Commit 4：Relay Pump 与第二跳 RDMA

[TransferEngine] Add the PXN relay pump

实现：

PxnPump 构造时创建一个独立 std::thread，持续推进 sender 和 relay，不占用调用线程或 RDMA WorkerPool thread 0。
Pump 固定执行 outbound completion、inbound child completion、inbound doorbell、outbound CE copy；无进展时等待50µs，sender enqueue可主动唤醒。
RelayPipeline round-robin 扫描7条 inbound lane，校验 Desc，并按 session解析 relay进程本地 SegmentID。
relay child batch持有自己的 TransferRequest vector，限定调用 RdmaTransport::submitTransferTask()。
MC_PXN_INFLIGHT_DEPTH 作为所有 lane共享的第二跳 Piece上限。
同 lane completion按 sequence连续发布，第二跳 RDMA 的64KiB切片、post和CQ poll全部复用既有 WorkerPool。
WorkerPool 零改动。

Commit 4 的调用关系为：

PxnPump std::thread
  ├─ SenderPipeline::reapOutbound()
  ├─ RelayPipeline::reapInbound()
  │    └─ child getTransferStatus() → Completion
  ├─ RelayPipeline::progressInbound()
  │    └─ doorbell → Desc → child RdmaTransport::submitTransferTask()
  │         └─ 既有 RDMA WorkerPool post/poll
  └─ SenderPipeline::progressOutbound()
       └─ ready event → CE copy → doorbell

本阶段新增一个 pxn_pump_test smoke case，覆盖 doorbell→第二跳提交→Completion。

Commit 5：接入 MultiTransport

[TransferEngine] Route eligible RDMA writes through PXN

实现：

Linux+CUDA 下，rdma factory在 MC_PXN_ENABLE=1 时构造 PxnRdmaTransport。
submitTransferTask() 识别 CUDA source WRITE；同轨、READ、CPU source或没有匹配 relay时限定调用基类。
根据 target buffer和Topology选出目标 HCA，再经 rail map标准化；目标 canonical rail与本地 rail不同时，从 Registry选择该 rail上的 relay。
可走PXN的 task按 target session聚合，再按8MiB/4096 plan切 Piece，不按传输总量或 Piece大小豁免。
enqueue首个 Piece前先建立全部 logical Slice和 task计数；Completion继续落到原 task的 success/failed/transferred_bytes。
sender enqueue后唤醒 PxnPump；公开 API、协议名和 getTransferStatus路径不变。
PXN禁用或初始化失败时仍使用普通 RdmaTransport。
API 和 decode 协议不变。

完整的 direct/eligible/mixed task list和真实 Registry选择放到 Commit 6健康路径E2E，不在本阶段继续扩充单测文件。

Commit 6：健康路径 E2E、benchmark 与运维文档

[TransferEngine] Add PXN end-to-end validation

覆盖：

同轨直发。
跨轨10KiB通过PXN relay。
8MiB、8MiB±1。
18 Piece恰好填满 ring。
第19个 Piece等待并复用 slot 0。
5GiB/640 Piece 多轮环复用及数据校验。
7 个 sender 同时写一个 relay 的7条 lane。
relay 无上层任务时主动转发。
健康路径必须满足 relayed_bytes > 0 && fallback_bytes == 0。

本阶段先证明正常传输链路、ring 复用和性能成立，不混入完整故障恢复实现。

Commit 7：Relay RDMA 提交快路径

[TransferEngine] Optimize the PXN relay RDMA submission path

目标：去掉 relay 第二跳在每个 Piece 上重复经过的通用传输对象和设备选择，保留现有 RDMA WorkerPool、QP、CQ 和背压实现。

实现：

RdmaRelayBackend 在 staging arena 完成 MR 注册后，一次解析并缓存 arena 对应的 RdmaContext*、本地 buffer范围和 lkey。Backend、Pump和 arena均先于基类 RDMA context销毁，因此该指针只在这段固定生命周期内使用。
target session第一次出现时建立 route cache。cache按 session保存本地 SegmentID，并按目标地址区间保存 {range, rkey, peer_nic_path}；后续 Piece只做区间命中，不再重复调用 getSegmentID()、getSegmentDescByID()或selectDevice()。
增加 PXN relay专用的 Slice构造入口。它直接根据 arena lkey和目标 route把一个 Piece切成现有 RDMA Slice，预填 source_addr、dest_addr、length、source_lkey、dest_rkey、peer_nic_path和 retry字段。
64KiB逻辑 Slice切分仍在该轻量构造入口完成，并且不得跨目标 buffer/rkey边界。WorkerPool本身不会把一个8MiB Slice再次切成64KiB，因此不能把整块 Piece直接压入队列。
给 RdmaContext/WorkerPool 增加接受 pre-resolved Slice的轻量提交入口，复用现有 prepared queue。该入口不再做 source/target metadata查找或设备选择，但继续使用现有 shard queue、endpoint建连、QP分摊、WR/CQ水位、CQ poll、retry和背压。
Pump不得直接调用 RdmaEndPoint::submitPostSend()，避免绕过 WorkerPool队列、连接状态和WR/CQ credit。
relay completion改为 Piece级轻量计数器，由所有子 Slice共享；最后一个 Slice完成时记录成功或失败。RelayTransfer::poll()直接读取该计数器，不再调用getTransferStatus()。
Slice从 RdmaRelayBackend持有的预分配 pool取得并在 Piece完成后归还。relay热路径不再构造 TransferRequest、BatchDesc、TransferTask，不再调用allocateBatchID()/freeBatchID()，也不做逐 Slice new/delete。
普通 RDMA提交路径和公开 Transport API保持不变；该入口只供 PXN relay内部调用。
cache miss只负责建立一次 route后进入快路径。metadata刷新、目标路由失效和重试时的cache重建统一放到 Commit 8，不在本 commit复制一套故障恢复状态机。

Commit 7 的调用关系为：

PxnPump std::thread
  └─ RelayPipeline::progressInbound()
       └─ RdmaRelayBackend::submit()
            ├─ arena route cache → RdmaContext* + source lkey
            ├─ session route cache → target_id + dest rkey + peer_nic_path
            ├─ 按 64KiB 和目标 MR 边界构造 pooled Slice
            └─ RdmaContext::submitPreparedPostSend()
                 └─ WorkerPool prepared queue
                      ├─ endpoint/QP/WR credit → ibv_post_send
                      └─ CQ poll → Piece completion counter

验收：

在现有 pxn_pump_test 增加一个快路径 case，不新增测试文件。
cache hit时对每个 Piece的 allocateBatchID、getSegmentDescByID和selectDevice调用次数均为0。
steady state relay提交不产生 TransferRequest/BatchDesc/TransferTask及逐 Slice堆分配。
Commit 6 benchmark保留为优化前基线；本 commit用相同负载复测10KiB、8MiB和5GiB，确认吞吐提升来自提交路径缩短，且WorkerPool queue、CQ poll和背压行为未改变。

4. 测试与验收

每个功能 commit 同时包含测试。

构建矩阵：

Linux、USE_CUDA=OFF
CUDA 12.8
CUDA 13
Debug + ASAN
CUDA/非 CUDA wheel smoke
实际 GPU/NVLink/RDMA

重点测试：

pxn_lifecycle_test
pxn_sender_test
pxn_pump_test
pxn_fault_test
pxn_transport_test
rdma_transport_submit_task_test
worker_pool_rail_state_test
ctrl_frame_test
sender_credit_test
transport_uint_test
graceful_shutdown_test

最终使用现有 SGLang disaggregation PP accuracy 测试分别验证 PXN 关闭和开启，不采用旧 Python PXN 实现。

性能门槛：

MC_PXN_ENABLE=0 时线程数和 HBM 与基线一致。
开启后每 rank 增加 kRequiredArenaSize（当前1008MiB），并在退出时完整释放。
同轨吞吐回退不超过3%。
健康跨轨吞吐达到同轨 native RDMA 的至少85%。
Pump 空闲 CPU 不超过单核5%。
无 CQ poll timeout 或明显 poll 间隔放大。
第一跳 CE 与前一 Piece 的第二跳 RDMA 能够重叠。
健康路径不存在逐 Piece stream/device synchronize。

5. 提交与代码风格

每阶段完成后执行：

git diff --check
pre-commit run --files <本阶段 changed files>

用户暂存后执行：

git diff --cached --check
pre-commit run --files $(git diff --cached --name-only --diff-filter=ACMR)

禁止 pre-commit run --all-files。

源码遵循 Mooncake 风格：

Google C++、4 空格、80 列、成员名后缀 _。
使用现有 Status、glog、RAII 和早返回。
注释仅解释 ownership/lifetime、ring 内存序、CUDA/RDMA 非直观约束。
不写逐行翻译式注释、阶段旁白、源码 ASCII 架构图、未来 v2 设想或注释掉的旧代码。
设计理由放在开发文档，行为由测试名表达。
每阶段由 Codex提供变更清单、测试结果和建议 commit message，用户确认后手工提交。

Leoda

Explorer

1.1. dev record