本文基于 HCCL e8351f6c 与 HCOMM ba1e77785 版本源码，深入分析 Ascend 950DT/A5 Host RDMA 通路是否具备多路径传输能力、依赖哪些散列输入实现，以及当前开源代码能确认到的工程边界。同时，全文系统梳理了 HostDPU 判定条件、执行流程、多 QP 分片机制、网卡硬件 LB 以及 UDP 源端口散列配置的协同逻辑。

核心结论

基于当前开源代码实现，可以明确以下五项核心事实：

1. Host RDMA 已接入实际数据面，算子覆盖有明确边界：950DT/A5 已经形成从 HostDPU 算法选择、HCOMM 资源申请，到 Host Endpoint、Channel、QP 建链、内存注册和 RDMA 读写的主通路实现，并非仅有抽象接口。目前已接入 AllReduce、AllGather、ReduceScatter、Reduce、Broadcast、Scatter、Barrier 等集合通信算子，AllToAll 家族（AllToAll/V/VC）以及点对点 Send/Recv/BatchSendRecv；但并非所有算子均已覆盖，例如 AllGatherV 和 ReduceScatterV 当前尚未实现 HostDPU 算法选择器。
2. Host RDMA 由物理拓扑自动选择，不存在全局独立总开关：CheckHostDPUOnly 判定仅检查 RankGraph 的最高通信层。该层必须存在覆盖通信域全部 Rank、包含 Host Endpoint 且严禁混入 Device Endpoint 的 CLOS 拓扑实例；链路通信协议还须为 RoCE，HCOMM 才会创建 HostCpuRoceChannel。仓内构造的测试用例涵盖单 SuperPod 两层拓扑（L1 使用 Host RoCE）与多 SuperPod 三层拓扑（L2 使用 Host RoCE、SuperPod 内 L1 使用 Device UB CLOS），真实部署通路以实际 RankGraph 组网为准。
3. 具备多 QP 消息分片与两类独立散列输入，但存在端口截断陷阱：
   • 每个 QP 可独立配置 RoCEv2 UDP 源端口，供网络交换机进行 ECMP 五元组散列；
   • 搭配 Hi1825/SP561 等 HRN 网卡时，可通过 HCCP 为每个 QP 下发独立的硬件散列输入 lbValue（即 hash_value）。
   • 关键工程陷阱：在默认 queueNum=1 下，即便网卡 LB 自动将 QP 数量扩展为 lbMax=N，由于源端口列表在建链配置阶段已被裁成 1 项，导致自动扩展出的 N 个 QP 全部退化为使用第 1 个 UDP 源端口。若要让 M 个 QP 真正各自绑定 M 个不同 UDP 源端口，必须显式配置 HCCL_RDMA_QPS_PER_CONNECTION=M（或通过 /etc/hcomm.cfg / IP 映射文件提供配置）。
4. “支持多路径散列”本质是流级静态 Hash，不保证物理路径分离与动态改路：UDP 源端口和 lbValue 仅向网络交换机或网卡内部提供散列因子。数据包在物理网络中的实际承载路径仍取决于网卡硬件、Provider 驱动映射、ECMP 策略和拓扑连线。当前 Host CPU RoCE 通路未启用 HCCP 头文件中定义的原生多路径（MPATH）与自适应路由（Adaptive Routing），亦未发现基于拥塞状态的动态重选路或 Host 多网卡主备切换逻辑。
5. QP 级异常检测与重传能力已接入，但 Channel 级自愈尚未闭环：代码已配置 QP 硬件超时重传，并实现 CQ Error 与 DFX 异常上报；但 HostCpuRoceChannel::Clean() 和 Resume() 当前为空实现（直接返回成功），表明 Channel 层尚未实现 QP 自动销毁重建或在线迁移。无法由硬件重传恢复的错误会向上返回，后续如何处置取决于上层通信域逻辑。

Host RDMA 通路、启用与执行

在探讨多路径前，首先需要理清 Host RDMA 在 950DT/A5 上的通路形态、自动触发逻辑与运行时执行闭环。

两种 Host RoCE 使用方式

在 CANN 通信体系中，需要严格区分以下两类 Host RoCE 通路：

本文重点分析第一类，即 HCCL 集合通信中 HostDPU 算法采用的 Host CPU + RoCE 通路。在此上下文中：

• Host 描述 Endpoint 位于主机侧；
• RoCE 描述网络通信协议；
• CPU 描述提交通信任务的执行引擎。三者不可混为一谈。

在实现上，CpuRoceEndpoint 仅接受 DEV_TYPE_950 与 DEV_TYPE_960 设备类型。950/960 设备在创建 HostRdmaConnection 时指定使用 OPBASE_QP_MODE，表明 950DT/A5 属于正式支持范围。

• CpuRoceEndpoint 实现
• HostRdmaConnection 实现
• HostCpuRoceChannel 实现

HCCP 在调用链路中的定位

HCCP（Huawei Collective Communication Protocol）并非“Device RoCE”的代名词，而是 HCOMM 内部承接 Socket、RDMA 资源管理和底层驱动调用的通信协议模块。Host RDMA 同样依赖 HCCP 完成 RDMA 设备初始化、MR 注册、CQ/QP 创建、QP 属性修改以及网卡 LB 扩展能力的调用。

HCCP 的封装层、Agent 与 Service 源码已随 HCOMM 开源（位于 src/base_comm/resources/hccp/）。然而，其在运行时通过 HccpDlopen 动态加载的 HRN Provider 专有动态库（如 libhrn5-rdmav34.so 和 libhrn5.so.1）并不在开源仓中。因此，开源代码能够精确确认参数如何组装并传递至 Provider 边界，但无法进一步确认 lbValue 在闭源网卡驱动及微码内部如何映射到具体物理队列或芯片出端口。

Host RDMA 如何被自动选中

Host RDMA 没有提供独立的全局显式总开关，而是由 HCCL 根据集群物理网络拓扑（RankGraph）自动识别并触发。

CheckHostDPUOnly 按以下规则逐级判定当前通信域是否属于 HostDPU 拓扑：

1. 多 Server 检查：topoInfo->serverNum > 1，单机场景直接退出 HostDPU；
2. 多层拓扑检查：topoInfo->topoLevelNums > 1，单层拓扑不进入 HostDPU；
3. 仅检查最高通信层：遍历层级时跳过 netLayer < (topoInfo->topoLevelNums - 1)；
4. CLOS 组网实例：最高层中查找 COMM_TOPO_CLOS 类型的拓扑实例；
5. 全 Rank 覆盖：候选 CLOS 实例覆盖的 Rank 数量必须等于通信域总卡数（rankNum == topoInfo->userRankSize）；
6. 无 Device 链路排除项：最高层拓扑实例中一旦发现任何 ENDPOINT_LOC_TYPE_DEVICE 的 Endpoint，立刻判定非 HostDPU 并退出；
7. Host 节点命中：至少发现一个 ENDPOINT_LOC_TYPE_HOST 的 Endpoint，将 hostDPUOnly 置为 true。

这段判定只关注最高层的 Endpoint 物理归属与 CLOS 覆盖完整性，本身不校验通信协议。一旦判定为 HostDPU，AutoSelector 立即配置执行引擎：

1
2

opParam.opExecuteConfig = OpExecuteConfig::HOSTCPU;
opParam.engine = CommEngine::COMM_ENGINE_CPU;

在后续资源申请阶段，HCCL 会针对具体链路按本地 Endpoint 类型进行拆分：Device Endpoint 使用 COMM_ENGINE_AICPU_TS，Host Endpoint 使用 COMM_ENGINE_CPU。当且仅当 Host Endpoint 的 Channel 协议为 RoCE 时，HCOMM 工厂方法才会实例化 HostCpuRoceChannel。因此，“通信域命中 HostDPU 算法”与“某层具体链路实例化为 Host RDMA 通道”是前后衔接但层级不同的两步。

• HostDPU 拓扑判定 CheckHostDPUOnly
• 自动选择器 AutoSelectorBase::Select
• HCOMM 通道工厂 ChannelFactory::CreateChannel

L0/L1/L2 分层拓扑与测试用例

为验证 HostDPU 代码路径，仓内在 LLT/ST 仿真环境（topo_model.cc）中构造了典型的分层组网模型。理解该模型有助于明晰 Host RoCE 在不同组网规模下的部署位次：

注意：拓扑层级为逻辑模型。 L0/L1/L2 是 RankGraph 中的逻辑网络层（NetLayer），并非固定绑定某个物理设备。在仓内单 SuperPod 两层测试用例中，最高层 L1 构造成 Host RoCE，因而 SuperPod 内跨 Server 通信直接走 Host RoCE；而在多 SuperPod 三层测试用例中，L1 是 Device 侧 UB CLOS，L2 才是 Host RoCE。生产环境中最终是否使用 Host RoCE，以部署时集群生成的实际 RankGraph 为准。

算子覆盖矩阵

通过对各算子 AutoSelector 中 SelectDPUAlgo 覆写情况的全面检索，当前算子支持矩阵如下：

由此可见，不能简单概括为“所有集合通信均支持 Host RDMA”。在当前版本中，AllGatherV 与 ReduceScatterV 的自动选择器尚未提供 HostDPU 选择分支。

HostDPU 任务执行流水线

HostDPU 的执行并非算法线程直接调用裸 verbs API，而是采用了“设备侧调度编排 + 共享内存控制桥 + Host 侧并发通信”的协同架构：

sequenceDiagram
autonumber
participant A as HCCL / AICPU编排
participant B as HCOMM任务桥
participant C as HostDPU模板
participant D as HostCpuRoceChannel
participant E as 对端Rank

rect rgb(239, 246, 255)
  Note over A,D: 阶段 1：资源准备
  A->>B: GetAlgResDPU：申请 100MiB DPUTAG 共享内存
  Note over B: 前半区 NPU→DPU 请求 / 后半区 DPU→NPU 响应
  A->>D: 按 Endpoint 创建 Channel<br/>Device→AICPU_TS, Host→CPU
end

rect rgb(245, 243, 255)
  Note over A,C: 阶段 2：任务控制路径
  A->>B: 序列化 DPURunInfo（模板名、参数、Channel、Rank、超时）
  A->>B: HcommSendRequest(npu2DpuShmemPtr)
  B->>C: HcclLaunchDPUKernel(request)
  activate C
  C->>C: 反序列化 → 查找模板 → DPUKernelRun
  par Device Endpoint
    A->>A: AICPU_TS 执行设备侧搬运
  and Host Endpoint
    C->>D: 调用 Host 侧 HCOMM 原语
  end
end

rect rgb(236, 253, 245)
  Note over C,E: 阶段 3：Host RDMA 数据路径
  E-->>D: NotifyRecord(NOTIFY_IDX_ACK)<br/>接收缓冲区就绪
  C->>D: NotifyWait(NOTIFY_IDX_ACK)
  C->>D: WriteWithNotify(NOTIFY_IDX_DATA_SIGNAL)
  D->>E: RDMA_WRITE_WITH_IMM<br/>数据写入远端 RMA Buffer 并携带 Notify ID
  C->>D: NotifyRecord(NOTIFY_IDX_FIN_ACK)
  C->>D: ChannelDrain
  C->>C: HcommFenceOnThread
  E->>E: NotifyWait(NOTIFY_IDX_DATA_SIGNAL)<br/>确认数据接收完成
  E->>E: NotifyWait(NOTIFY_IDX_FIN_ACK)<br/>确认发送端排空
end

rect rgb(255, 247, 237)
  Note over A,C: 阶段 4：完成收口
  C-->>B: 模板返回执行状态
  deactivate C
  B-->>A: HcommWaitResponse(dpu2NpuShmemPtr)
  A->>A: 校验 msgId 并继续后处理
end

1. 共享内存分账：HCCL 在 GetAlgResDPU 中申请 100MiB 且打上 DPUTAG 标记的设备共享内存，以 DPU2NPU_SHMEM_RATIO = 2 平分为前半区（NPU→DPU 请求）和后半区（DPU→NPU 响应）。
2. 任务描述符下发：算法将执行上下文 DPURunInfo（包含模板名、切分参数、Channel 句柄、Rank 与超时配置）序列化写入请求区，通过 HcommSendRequest 投递至 Host 侧。
3. 模板派发与协同：Host 侧 HcclLaunchDPUKernel 反序列化请求，经 InsAlgTemplateRegistry 获取对应模板实例并调用 DPUKernelRun。此时 Device 侧可并行处理本节点 UB 内存搬运，Host 侧则由 Host CPU 提交通信。
4. 三段式握手与数据传输（普通 SendRecvWrite 路径）：
   • Step Sync（就绪确认）：接收端调用 NotifyRecord(NOTIFY_IDX_ACK)，发送端调用 NotifyWait(NOTIFY_IDX_ACK)，确保远端目标内存就绪；
   • 数据写入：发送端调用 HcommWriteWithNotifyNbiOnThread 下发数据写，尾块携带 NOTIFY_IDX_DATA_SIGNAL；
   • Fin Sync 与 Drain（完成排空）：发送端触发 NOTIFY_IDX_FIN_ACK 并调用 ChannelDrain 与线程 Fence；接收端先后等待 NOTIFY_IDX_DATA_SIGNAL 与 NOTIFY_IDX_FIN_ACK 确认数据完全可见。
5. 任务收口：模板执行结束后向响应区写入状态，AICPU 侧通过 HcommWaitResponse 校验 msgId 并结束算子调度。

上述时序图展示的是普通 SendRecvWrite 包装器。部分 NHR 算法使用 DpuBatchTransfer 优化路径，将 STEP_SYNC、批量 WriteWithNotify、Drain/Fence 与 DATA_SIGNAL 等待组合执行，不使用 NOTIFY_IDX_FIN_ACK；因此不能把图中的 FIN_ACK 流程视为所有 HostDPU 算法的统一时序。

控制路径与业务数据路径分离。 HcommSendRequest/HcommWaitResponse 在 100MiB DPUTAG 共享内存中传递任务元数据；Host RDMA 阶段的业务载荷则通过已注册的 RMA Buffer 与 RDMA QP 传输，不经 DPUTAG 任务共享内存中转。

• HostDPU 共享内存资源分配
• HostDPU 任务派发 HcclLaunchDPUKernel
• HostDPU 数据传输包装器 SendRecvWrite

异步建链状态机

HostCpuRoceChannel 采用异步状态机设计，每次调用 Connect() 驱动当前阶段向前推进：

stateDiagram-v2
direction TB
[*] --> INIT
INIT --> SOCKET_OK: CheckSocketStatus()<br/>建立或复用辅助 Socket（默认端口 60001）
SOCKET_OK --> CAP_EXCHANGED: ExchangeCapability()<br/>交换版本、通信魔数与特性能力
CAP_EXCHANGED --> QP_CREATED: CreateQp()<br/>按 queueNum 或网卡 LB 能力创建 CQ/QP
QP_CREATED --> DATA_EXCHANGE: ExchangeData()<br/>交换 QPN、PSN、GID、MR 地址/RKey 等
DATA_EXCHANGE --> QP_MODIFIED: ModifyQp()<br/>配置 SL、TC、重传、源端口与 lbValue
QP_MODIFIED --> READY: SyncAfterModifyQp()<br/>每个 QP 预投递 16 个 Receive WQE
READY --> [*]

带外 TCP Socket 仅用于能力协商、元数据交换与状态同步，通道进入 READY 状态后，所有业务数据与通知均在 RDMA QP 上传输。

内存注册与 D2N=UB 代理机制

Host RDMA 读写的源地址和目的地址必须严格落在已注册并交换的 RMA Buffer 范围内，否则 FindLocalBuffer 或 FindRemoteBuffer 会返回 HCCL_E_NOT_FOUND。

CpuRoceEndpoint 以 {devPhyId, COMM_PROTOCOL_ROCE, ENDPOINT_LOC_TYPE_HOST, ipAddr} 作为唯一键缓存复用底层 RDMA 上下文，本地 MR 亦在此粒度进行管理。在 950DT/Hi1825 架构下，Device 内存注册存在一个特殊设计：

• 普通内存注册：Host 内存及常规 Device 内存直接按用户传入的连续虚拟地址区间 [va, va+size] 进行注册。
• D2N=UB 别名注册策略：当 Host 网卡与 NPU 之间的 Device-to-NIC（D2N）数据通路采用 UB 总线时，底层 ubdevshm 驱动如果对重叠 VA 进行重复登记会直接返回错误（如 -17）。由于当前 HCCP 缺乏直接探测 D2N 物理总线类型的接口，代码暂以 capabilities_.lbMax > 0 作为命中 Hi1825/UB 架构的代理条件。
• 对满足条件的 Device 内存，RoceRegedMemMgr 会调用 aclrtMemGetAddressRange 获取底层整段 aclrtMalloc 分配的大内存块，针对整段 allocation 注册物理硬件 MR 并建立引用计数；用户传入的实际数据子区间仅作为共享同一物理 LKey/RKey 的 Alias Buffer 导出。

代码注释已明确标明：该代理逻辑是当前版本的过渡设计，未来 HCCP 提供专有 D2N/UB 探测接口后将予以替换。

• RoceRegedMemMgr 实现与 D2N=UB 代理注记

数据操作与通知语义

针对标准对称 Host CPU RoCE 模式，通道实现了完备的原语支持：

1. Write 与 Read 分片调度

CpuRoceEndpoint 目前将单次 RDMA 操作上限暂定为 1GiB（RDMA_MAX_WR_LENGTH = 1GB）。数据发送时，首先按 1GiB 切割成若干 Chunk，每个 Chunk 再根据参与传输的 QP 数量执行 Striping 分片。多 QP 开启时，HCCL_MULTI_QP_THRESHOLD 决定了实际承载有效载荷的 QP 个数：

当初始平均切片长度满足 tileLen != 0 && tileLen < qpThreshold 时，代码按下式缩减实际承载数据的 QP 数量：

$$useQpNum = \min\left(actualQpNum,\ \left\lfloor\frac{tailLen - 1}{qpThreshold}\right\rfloor + 1\right)$$

若 tileLen == 0（即数据长度小于 QP 数量），该缩减逻辑不会进入，useQpNum 保持为 actualQpNum；此时前面的 QP 获得 0 长度切片，尾 QP 承载剩余数据。

对于未分配到有效载荷的剩余空闲 QP（$i \ge useQpNum$），仍会提交长度为 0 的 Signaled WR。此设计至关重要：由于 WaitForWqeCompletion() 需要轮询等待通道内所有 QP 的 wqeNums_[i] 计数清零，为所有 QP 统一投递 Signaled WR 能确保多 QP 的完成状态推进严格保持步调一致，避免特定 QP 挂起引发死锁。

2. NotifyRecord 与 NotifyWait

• NotifyRecord：向通道内所有 QP 提交一条长度为 0、携带 32 位 Immediate Data（存放目标 Notify ID）的 RDMA 写操作；
• NotifyWait：遍历通道内所有 QP 对应的 recvCq，逐一轮询并校验出队的 wc.imm_data == dpuNotifyId；
• 滑动信标补齐：单次 NotifyWait 完成后，通道会立即调用 IbvPostRecv() 为每个 QP 各自补投递 1 个 Receive WQE，使 QP 的接收缓冲池始终维持在 16 个信标深度；
• 若轮询超时，直接返回 HCCL_E_TIMEOUT；若 CQE 返回异常，则提取本地/远端 Server、Device ID、IP 及 QP 状态并上报 DFX。

3. WriteWithNotify

大于 1GiB 的前置数据块采用普通 RDMA_WRITE 发送；最后一个尾块在所有 QP 上调用 RDMA_WRITE_WITH_IMM。承载数据的 QP 携带切片数据并打上 Imm Notify，空闲 QP 发送长度为 0 的 Imm Notify，保证接收端每个 QP 都能收到一致的同步信号。

4. ChannelDrain 与 ChannelFence

• 两者均通过 WaitForWqeCompletion() 轮询等待所有 QP 的所有在途 Send WQE 完全出队清空；
• ChannelFence 在排空后将 fenceFlag_ 置为 true，使得后续第一条投递的 RDMA WR 自动附带 IBV_SEND_FENCE 标志；
• HCCL HostDPU 传输包装器优先调用轻量级的 Channel Drain，并在其后叠加线程级 HcommFenceOnThread。为向后兼容低于 CANN 9.2.0.2 的老版本，HCCL 会自动将 Host Channel Drain 回退映射为 Channel Fence。

多QP、LB 与多路径能力

本节是全文的核心：厘清数据分片多 QP、网卡硬件 LB 与 RoCEv2 UDP 源端口三者的机制定位，进而评估其多路径能力与边界。

QP 数量解析与优先级

HCCL_RDMA_QPS_PER_CONNECTION 定义了 Rank 间单个逻辑 Channel 期望分配的 QP 数量（合法范围 1～32，默认 1，推荐 $\le 8$）。HCCL_MULTI_QP_THRESHOLD 定义了每个数据 QP 的最小切片阈值（范围 1～8192KB，默认 512KB）。

在 A5 CommunicatorV2 框架下，Channel 的 queueNum 按照以下绝对优先级逐级解析（见 ResolveQueueNum）：

flowchart TD
A{"调用方是否显式指定<br/>HcclChannelDesc.roceAttr.queueNum？"}
B["采用调用方配置"]
C{"Host RoCE 的 /etc/hcomm.cfg<br/>是否配置 QP 数？"}
D["采用 hostMultiQpConfig.GetQpCount(devicePhyId)"]
E{"本端/远端 IP 对是否命中<br/>MultiQpSrcPort.cfg？"}
F["采用 IP 对对应的端口数量"]
G["采用 HCCL_RDMA_QPS_PER_CONNECTION<br/>未配置时默认为 1"]
A -- 是 --> B
A -- 否 --> C
C -- 是 --> D
C -- 否 --> E
E -- 是 --> F
E -- 否 --> G

建链对称性约束： 建链阶段通信双方会严格校验对端声明的 QP 数量。若通信双方环境变量不一致，或两端网卡硬件 LB 能力差异导致扩展出的实际 Connection 数量不同，建链会报错退出，不会自动协商取较小值降级运行。

网卡硬件 LB

950DT/A5 搭配 Hi1825/SP561 网卡时的 LB 交互时序如下：

1. 能力探测：CpuRoceEndpoint::GetCapabilities 调用 HCCP 的 RaGetLbMax；
2. Provider 符号解析：HCCP 动态绑定 HRN Provider 动态库并调用 roce_get_qp_num。若符号不存在，优雅回退为 lbMax = 0；
3. 分配哈希因子：HostCpuRoceChannel::BuildConnection 遍历各连接，为第 $i$ 个 QP 赋值 qpInfo.lbValue = i % lbMax；
4. 属性下发：QP 推进至 RTS 状态后，HostRdmaConnection::ModifyQp 显式调用 RaSetQpLbValue(qpHandle, lbValue)，最终由 Provider 的 roce_set_qp_lb_value 写入网卡。

根据配置 QP 数与网卡 lbMax 的不同组合，实际 QP 数量与 lbValue 的对应关系如下：

• HostCpuRoceChannel 的 LB 与多 QP 分配逻辑
• HCCP LB 接口声明
• HCCP 对 HRN Provider 扩展符号的动态绑定

Host RDMA UDP 源端口与 ECMP 散列

950 系列支持通过环境变量 HCCL_HOST_RDMA_UDP_PORTS_LIST 按 NPU 物理设备 ID 为 Host RoCE QP 分配指定的 UDP 源端口：

1
2

# 格式：<phy_dev_id>:<port1>,<port2>...;<phy_dev_id>:<port1>,<port2>...
export HCCL_HOST_RDMA_UDP_PORTS_LIST="0:10000,10015;1:10016,10031"

配置必须遵循以下规则：

• phy_dev_id 须为非负十进制整数，同一物理设备 ID 严禁重复；
• 端口号必须位于 $[1, 65535]$，应避开系统保留端口 $[1, 1023]$；
• 单个设备端口数量上限为 32 个，环境变量总字符长度上限为 32KiB；
• 格式错误会导致通信域初始化异常退出。

1. queueNum、LB 与源端口协同组合（关键陷阱揭秘）

在 HCOMM 实现中，“配置阶段的 QP 数”与“Channel 内部实际建立的 QP 数”之间存在处理时序差：

1. Step 1：HCOMM 首先解析出 Channel 描述符中的 channelDesc.roceAttr.queueNum（若未显式指定，默认为 1）；
2. Step 2：RoceChannelDescConfigurator::FillRoceSrcPortList 根据上述 queueNum 调整并裁剪端口缓存：若配置的端口过多则截断，过少则循环复用，最终缓存数组大小严格等于 queueNum；
3. Step 3：进入 HostCpuRoceChannel::BuildConnection，查询网卡能力 lbMax。仅当 queueNum == 1 && lbMax > 0 时，创建循环次数才从 1 自动扩展为 lbMax；
4. Step 4：循环内部为第 $i$ 个 QP 赋值源端口：

   1	qpInfo.udpSport = srcPortList[i % channelDesc_.roceAttr.queueNum];

这形成了一个需要特别注意的实现组合约束：

配置建议： 如果希望建立 $M$ 个 QP 并使其分别获得不同的 UDP 源端口，不能仅依赖网卡 LB 自动扩展。需要将 queueNum 显式解析为 $M$（例如配置 HCCL_RDMA_QPS_PER_CONNECTION=M，或通过 /etc/hcomm.cfg、IP 对配置得到同一数量），并保证提供至少 $M$ 个互不相同的端口。

2. ECMP 五元组哈希机制

RoCEv2 数据包在 IP 层的五元组为：源 IP、目的 IP、源 UDP 端口、目的 UDP 端口（固定 4791）、IP 协议号（固定 17 UDP）。在固定的 Rank 对之间通信时，两端 IP、目的端口及协议号均不可变，因此修改源 UDP 端口是主机端为交换机 ECMP 提供散列输入的最主要途径。

但必须指出：源端口不同仅向交换机哈希桶提供了离散输入；各 QP 最终是否会落到不同交换机物理链路上，取决于交换机具体的逐流哈希算法、网络拓扑等宽情况以及哈希碰撞几率。

3. 源端口的最终覆盖层级

源端口的生效遵循“环境解析”与“通道绑定”两级映射：

1. 设备级端口基线：首先解析 HCCL_HOST_RDMA_UDP_PORTS_LIST；若本地存在 /etc/hcomm.cfg，其中配置的同物理设备项（须同时提供 udp_port_mode_<id>=multi_qp、multi_qp_count_<id> 及 multi_qp_udp_ports_<id>）会覆盖该物理设备对应的环境变量配置；
2. Channel 级分配：创建具体 Channel 时，首先使用上述物理设备的端口列表；若未获取到，再按本地/远端 IP 对查询 MultiQpSrcPort.cfg；若仍未命中，向底层传递 udpSport = 0，由 Provider 驱动自行填充。

• HCCL_HOST_RDMA_UDP_PORTS_LIST 用户指南
• HCCL_RDMA_QPS_PER_CONNECTION 用户指南
• HCCL_MULTI_QP_THRESHOLD 用户指南
• HostMultiQpConfig 源码实现
• RoceChannelDescConfigurator 端口绑定源码

多路径能力边界与最终判定

综合代码剖析，Ascend 950DT/A5 Host RDMA 在多路径传输维度能够确认的能力边界如下：

最终结论：
Ascend 950DT/A5 Host RDMA 已经具备多 QP 数据分片能力，并可通过 UDP 源端口与网卡 lbValue 提供静态路径散列因子；在支持 ECMP 的网络架构中，这些因子可使多 QP 流量具备分散到多条网络路径的条件。该能力属于流级（QP 级）静态散列选路，不能保证各 QP 必然落到不同物理路径。从当前开源代码看，Host CPU RoCE 通路也未启用网卡原生 MPATH、自适应路由、动态拥塞感知选路或 Host 多网卡主备故障切换。

关联机制与工程边界

Host NIC 与 Device NIC 异构混合模式

在 HostCpuRoceChannel::Connect() 协商阶段，双方会交换通信栈能力。如果对端声明了 COMM_STACK_TRANSPORT_IBVERBS，通道将转入 Host NIC — Device NIC 混合模式（源码中命名为 Hybird）。Send/Recv 选择器中的 DpuSendSoleHost 与 DpuRecvSoleHost 即为此类跨节点异构链路的使用者。

混合模式与标准对称 Host CPU RoCE 存在本质区别：

• 单 QP 传输：退化为仅使用第 0 个 Connection/QP；
• 内存模型差异：重新注册用户 Buffer，并额外交换 2 个数据 Buffer 与 3 个 4 字节的 Host Notify 内存；
• 链式通知协议：WriteWithNotify 不使用 RDMA_WRITE_WITH_IMM，而是由一条数据 RDMA Write 串联一条写入对端 Notify 内存的链式 RDMA Write 组成；
• 内存原子轮询：NotifyWait 通过在本地原子轮询特定 Notify 内存字实现，默认超时 30 秒；完全不依赖标准模式下的 Receive WQE 和 recv CQ 机制。

这是一条专为兼容旧版 NPU RoCE/HCCP 设备端通信栈而设计的异构降级链路，其实现细节不代表 950 对称 HostDPU 主通路。

单边通信与 HIXL 模式

当 HcommChannelDesc.exchangeAllMems = true 时，通道进入单边通信模式（注释标为 HIXL 使用；HCCL 集合通信固定为 false）。

该模式跳过了常规 Channel 的逐个内存句柄传递流程，改为直接交换 Endpoint 上的全量已注册内存；非 Client 端会按需拉起监听 Socket。特别需要注意的是：在单边模式下，自动源端口分配机制被强制旁路，udpSport 固定下发为 0。因此在评估集合通信 Host RDMA 时，不应混入单边通信的参数规则。

QoS、重传与异常处理边界

Host RDMA 在 QP 属性修改阶段统一注入 RoCE 网络质量参数：

运行时具备以下异常保护与 DFX 诊断：

• verbs 提交阶段细分返回码：区分队列溢出（ENOMEM / HCCL_E_AGAIN）与网络错误（HCCL_E_NETWORK）；
• send/recv CQ 轮询检测 ibv_wc_status、Opcode 及 Vendor Error；
• NotifyWait、Drain 和 Fence 操作均受执行超时约束；
• CQ 发生异常时，自动收集本地/远端 Server、Device ID、IP 以及 QP 状态，并触发 DFX 任务结束诊断。

容错恢复的局限性：
HostCpuRoceChannel::Clean() 和 Resume() 目前直接返回 HCCL_SUCCESS。这意味着在当前版本中，一旦通信通道遭遇无法依靠 QP 硬件重传恢复的故障，HCOMM 无法在 Channel 级别自动重建 QP 或无缝重哈希改路；错误会向上返回，后续是否重建通信域由上层恢复策略决定。

NDA 设备直驱模式对比

除 Host CPU 驱动外，950 还支持 AICPU_TS 或 AIV 直接下发任务给 Host RoCE 网卡（即 NDA 通路）。该场景下本地 Endpoint 依然是 CpuRoceEndpoint（表示网卡位于 Host 侧），但 Channel 变为 AicpuTsRoceChannelV2。

由 DevRdmaConnectionV2 查询底层 RaNdaGetDirectFlag 并区分 DMA 模式：

• DIRECT_FLAG_PCIE $\to$ QBUF_DMA_MODE_DEFAULT（PCIe 总线模式）；
• DIRECT_FLAG_UB $\to$ QBUF_DMA_MODE_INDEP_UB（UB 总线模式）。

值得注意的是，当前设备侧 RdmaConnLiteV2 明确写明：AICPU NDA 暂不支持 PCIe 模式，仅在 UB 模式下加载 Rdma1825Ops；而 AIV 则支持 PCIe NDA。因此在定位现场问题时，必须结合 Endpoint 位置、Channel 协议、CommEngine 与 NDA DirectFlag 四个维度综合判定。

• DevRdmaConnectionV2 实现
• HcommChannelCreate 的 NDA 使用约束

生产部署排查核对清单

在生产集群中，若要确认业务流量是否真正进入 Host RDMA 以及多路径是否按预期生效，可依据以下步骤逐步核验：

1. 核对拓扑命中：查看初始化日志中 [CheckHostDPUOnly]，确认最高通信层是否为全覆盖、纯 Host Endpoint 的 CLOS 拓扑，无 Device Endpoint 干扰；
2. 核对算法与通道：确认选中的算子算法名为 Dpu*（如 DpuAllReduceSequenceMeshNHR），Channel 实例确认进入 HostCpuRoceChannel；
3. 核对 QP 数量来源：检查 queueNum 日志，确认其最终解析自显式 Channel 配置、/etc/hcomm.cfg、MultiQpSrcPort.cfg 还是环境变量 HCCL_RDMA_QPS_PER_CONNECTION；
4. 核对网卡 LB 能力：查看 RaGetLbMax 日志返回值是否 $> 0$；检查 HostCpuRoceChannel::BuildConnection 日志中每个 QP 的 lbValue（0 至 $lbMax-1$）；
5. 核对 UDP 源端口分布：检查每 QP 的 udpSport 日志输出。重点警惕：若发现所有 QP 的源端口完全相同，检查是否误用了“默认 queueNum=1 + 网卡自动扩展 LB”的组合；如需离散端口，须显式配置 HCCL_RDMA_QPS_PER_CONNECTION；
6. 核对交换机 ECMP 配置：确认物理交换机针对 RoCEv2（UDP 4791）开启了基于五元组（特别是源 UDP 端口）的哈希算法；
7. 核对物理网络流量：在业务运行期间采样网络交换机上联端口的流量计数器（PFC / SNMP / 遥测计数）。单靠 HCCL 日志显示多 QP 仅能证明软件层已下发散列因子，物理层是否真正负载均衡须以交换机物理端口流量分布为准。

导读

核心观点

后训练 RL 的竞争点正在从 换哪种训练目标（objective） 迁移到 能否稳定组织一条可观测、可回放、可扩展的训练链路。

本文结构与要点

全文围绕五条主线展开：系统视角、Runtime 系统、框架坐标、补充展开与落地选型。

除 objective 外，框架还需组织 rollout、reward、环境、serving、trace 回流与故障恢复。各部分核心论点如下：

系统视角

• 后训练 RL 已从算法问题变成系统工程问题
• 优化对象从单轮回答扩展到真实任务过程
• Reward 是系统接口，不仅是单标量分数
• KL 只能约束分布距离，无法约束任务语义
• RL 数据是动态轨迹资产，不是静态 dataset

Runtime 系统

• RL 框架正在从 trainer 变成 runtime
• RL 系统走向服务化与调度化，因为 trainer 抽象不够
• 生产态推理集群可作为低优先级 rollout 后端，而非简单混入线上请求路径
• OPD 将 teacher feedback 变成需要调度、观测和回放的高成本服务
• 大规模后训练的瓶颈常常在 rollout

框架坐标

• 框架按调度层、训练栈、代码谱系等维度分类
• 另有研究向、轻量实验、零侵入 Agent、fork 衍生等补充类别
• 社区路标集中在 multi-turn、async、tool、production

补充展开

• Objective / feedback 路线从 PPO、GRPO 扩展到 OPD
• Agentic RL 放大 credit assignment、token 对齐与稳定性问题
• 系统扩展规律从单点吞吐转向 loop goodput（有效吞吐）

落地

• 未来系统会分成应用、环境、编排、模型 runtime 四层
• 框架选型应围绕当前瓶颈做检查
• 竞争壁垒在训练链路，不在单个 trainer repo

文中会反复提到 DPO、GRPO、OPD 等方法名；各 objective / feedback 路线的形式对照见后训练 Objective 速查。

系统视角：从算法到系统工程

核心变化：从 RLHF 到 Agentic RL

观点：后训练 RL 的核心变化，不是 objective 公式突然颠覆，而是优化对象从“单轮回答偏好”扩展到了“真实任务过程”。

分析线索：对比旧 RLHF 链路与 agentic RL 训练链路。

早期 RLHF 常可概括为：

1

SFT -> Reward Model -> PPO

该链路仍成立，但已不足以描述当前的 post-training。更完整的流程可表示为：

1
2

任务 / 环境 / 工具 -> 轨迹采样 -> 可验证反馈 / judge / reward shaping
             -> policy update -> 新 policy 继续进入环境

对于 chat alignment，样本通常是 prompt 和 response，reward 通常反映对最终回答的整体偏好。在 agentic post-training 中，样本变为多轮轨迹：模型需读文件、调用工具、执行命令、观察结果、修正计划并提交 artifact。此时 reward 除最终回答质量外，还需覆盖过程可靠性、工具使用有效性与结果可验证性。

后训练 RL 的一个要点是：环境设计与算法选择同等重要。

弱环境配合强算法，可能导致模型 exploit reward 设计缺陷；清晰环境配合可验证信号，即便算法相对简单，也能提供稳定的优化方向。

Reward：系统接口，不仅是分数

观点：Reward 在工程实现中更接近跨系统接口，而非 trainer 直接消费的单标量。

分析线索：reward contract 要同时表达任务完成、质量、约束、过程、安全、成本与稳定性。

常见讨论将 reward 视为标量分数；在工程实现中，reward 更接近跨系统接口。

这个接口至少包含几层信息：

• 任务是否完成；
• 完成质量如何；
• 有没有违反约束；
• 过程里是否出现不可接受行为；
• 哪些子步骤值得鼓励或惩罚；
• reward 是否存在被模型 exploit 的风险；
• reward 的延迟、成本和稳定性如何。

在数学题、代码题、工具任务中，可验证 reward 可将部分主观偏好转化为客观检查：测试是否通过、命令是否成功、artifact 是否符合格式、答案是否可被程序验证。

但可验证并不等于简单。真实任务经常会出现这些问题：

• 单个最终结果可验证，但中间过程不可见；
• 测试集太窄，模型学会过拟合测试；
• judge 本身不稳定，导致 reward 噪声放大；
• reward 只关心结果，模型学会高风险路径；
• reward 过度 shaping 时，模型可能优化 shaping 信号而非任务本身。

因此 reward 设计类似 API 设计：接口语义不清时，trainer、rollout engine 与数据平台需承接由此产生的歧义与耦合。

可预期的是，post-training 团队将逐步沉淀各领域的 reward contract。coding、terminal、browser、office、search、math、data analysis 等场景各有不同的可验证条件、失败类型与 trace schema。通用 RL 框架主要提供集成接口，效果取决于具体领域 contract 的质量。

KL：分布约束，不是安全边界

观点：KL 是控制 policy drift 的手段，不是 agentic RL 的安全边界本身。

分析线索：区分 KL 能管什么，以及必须由 reward 反作弊、环境隔离、replay / eval 处理什么。

RLHF 中常见做法是通过 KL penalty 将 policy 约束在 reference model 附近，以限制 policy drift。在 PPO-RLHF 中，reward 提供优化方向，reference 与 KL 提供分布约束。

但在 agentic RL 中，KL 的作用需单独评估。

一方面，KL 仍可用于防止 policy 为短期 reward 牺牲语言质量、格式规范与基本行为边界。另一方面，agent 任务通常要求 policy 偏离 SFT 初始分布：更长规划、更主动的工具使用、更多自我修正与失败恢复。KL 过强时，policy 可能被约束在原有行为分布附近，难以学习新策略。

因此 KL 并非安全机制本身，而是控制 policy drift 的手段。稳定性还依赖：

• reference / KL；
• reward 反作弊设计；
• 环境权限隔离；
• rollout 过程审计；
• bad case replay；
• 离线 eval 与在线 shadow eval；
• 对工具、文件、网络、执行命令的边界控制。

概括而言：KL 约束分布距离，不约束任务语义。

RL 数据：动态轨迹资产

观点：RL 数据是随 policy 变化不断生成的 trajectory asset，而不是一次性准备好的静态 dataset。

分析线索：trajectory schema 至少要覆盖复现、归因和 replay 所需字段。

SFT 消费静态语料；RL 则持续生成新数据。

每轮 policy update 后，模型分布变化，下一轮 rollout 的错误类型、工具调用模式、回答长度与失败方式也随之变化。数据并非固定存储中的静态文件，而是由当前 policy 与环境交互动态产生。

这带来以下工程后果：

• 数据质量与 policy 版本强绑定；
• trace 需记录足够上下文以支持复现；
• reward 版本变化会影响跨轮比较；
• old policy / new policy / reference policy 的 logprob 须严格区分；
• 失败样本通常具有较高分析价值，可用于后续改进。

因此 RL 数据平台不宜仅按 dataset 模型设计，而更接近 experiment log、trajectory store、replay buffer 与 evaluation archive 的组合。

对框架来说，这意味着数据结构里至少要能表达：

• prompt / observation / action / response；
• token logprob；
• reward breakdown；
• judge version；
• environment version；
• tool call result；
• timeout / exception；
• policy checkpoint；
• sampling params；
• trace span。

缺少上述字段时，后续归因分析将十分困难。

Runtime 系统：从 Trainer 到 Runtime

Runtime 角色：组件与边界

观点：现在看 RL 框架，重点不是内置了几种 post-training 方法，而是能否组织一整条训练链路。

分析线索：框架抽象要同时解释 policy、reference、reward、rollout、environment、trainer、scheduler、data store、observability 的关系。

评估 RL 框架时，不应仅关注是否内置 PPO / GRPO 模块，而需评估其组织完整训练链路的能力。

一个 LLM RL 框架通常至少要处理这些角色：

这里最关键的是 policy update 和 rollout generation 的关系。

传统深度学习训练中，DataLoader 按 batch 加载样本，trainer 执行 forward/backward，数据路径与训练路径边界较为清晰。LLM RL 中，rollout generation 本身是一个分布式推理系统，且需频繁同步最新 policy，训练系统因而成为训推混合系统。

因此许多框架强调 Ray、vLLM、colocated placement、offload、async rollout、hybrid engine、placement group 等能力。这些并非“外围工程”，而是 RL 框架的核心组成部分。

服务与调度视角：Runtime 的系统边界

图示：RL 后训练系统架构图

观点：现代 RL 后训练系统的边界正在从“计算 loss 的 trainer”扩大到“组织跨服务 data loop 的 runtime”。

分析线索：将 rollout、reward、environment、policy publish / weight sync、trajectory store、observability 视为独立服务，分析调度层如何维持 goodput 与版本语义。

这里的 goodput 指有效吞吐：不是系统总共生成了多少 token 或处理了多少请求，而是单位时间内有多少可用、可验、版本语义正确的 trajectory 最终进入 reward 与 policy update。

RL 系统正在服务化

图示：从 Trainer 到 Runtime Loop

后训练 RL 系统不再只是一个 trainer 脚本加几段 rollout 代码，而是由多个职责清晰的服务协作：

这些服务之间的数据流构成 training loop：policy 更新触发权重同步或版本发布，Rollout Service 按新 policy 生成 trajectory，Environment Service 提供交互状态，Reward / Judge Service 产出可版本化的反馈，Trajectory Store 持久化全链路字段，Observability Service 将失败与长尾归因至具体环节。

核心变化是：RL 框架的边界从“计算 loss”扩大到“组织服务间的 data loop”。Rollout Service 承担 policy 的 generation forward（推理采样）；Trainer 承担 policy update（通常需在已收集 batch 上重算 logprob / value，再 backward 更新权重）。

系统能否扩展，取决于各服务能否独立扩缩、版本对齐与失败隔离。

OPD 亦会在 RL 系统中引入高成本的 Teacher / Distillation Service。传统 RLVR 的瓶颈在 rollout 与 verifier；OPD 还需在学生 trajectory 上调用 teacher，并将 token / chunk / step 级反馈回传 trainer。

服务接口因此需支持 teacher model version、teacher logits 或 compressed feedback、chunk-level preference / semantic score、feedback span 与 token span 对齐、teacher-student tokenizer mapping、feedback latency / queue delay。若 rollout 使用生产态推理集群的低谷容量，OPD 进一步要求调度层同时安排 student rollout 与 teacher feedback，并隔离二者的优先级、版本和队列延迟，RL 后训练系统因而趋向 multi-model inference orchestration。

Rollout Service 是训练系统里的在线推理服务

Rollout 无法以 trainer 内单次 model.generate() 概括。其本质是面向 RL 负载的在线推理服务，具有 serving 系统的典型特征：

• batching 与 continuous batching；
• KV cache 管理与 prefix 复用；
• speculative decoding；
• 长尾请求与 straggler；
• 权重热更新 / policy version publish 与 worker health check；
• 请求超时、局部重试与 worker 故障恢复。

但它又不同于面向用户的线上 serving：优化目标不是单请求 latency，而是 RL step 的 goodput——单位 wall-clock 内有多少有效 trajectory 进入 reward 和 update。这意味着 batch 策略、截断策略、并发上限都要围绕“样本吞吐”而不是“p99 延迟 SLA”来设计。

推理后端的选择反映 workload 差异。SGLang 侧重多轮结构化执行、KV 复用与程序式 rollout；vLLM 侧重通用 serving 与生态覆盖。框架侧大致为：slime / Miles / Relax 倾向 SGLang，OpenRLHF / vime 倾向 vLLM，verl / ROLL / NeMo-RL 支持双后端。选型应基于 workload 是单轮大批量、长上下文多轮，还是工具环境驱动的 agentic generation，而非单一 benchmark 分数。

生产推理集群可能成为 rollout 后端

图示：Rollout 后端部署形态

图示从左到右为三种演进形态：Colocated（rollout 与 trainer 同 job 部署，也称 embedded）、Disaggregated Rollout Service（为本训练任务单独部署 rollout 集群，经 queue / RPC 与 trainer 解耦）、Production Inference Rollout（训练过程中接入生产态推理平台，用低优先级 / 可抢占任务消纳低谷或弹性余量）。

后两种的数据流相近：trainer 下发 prompt，rollout 回传 trace + logprob；差别在于 2 通常是 trainer 将更新后的权重同步到训练专用 rollout worker，3 更像 trainer 将 policy version 发布到 serving 控制面，由生产态推理池按版本装载或路由。

传统 RL 框架默认 rollout engine 随训练 job 启动：训练集群同时承担 generation forward 与 policy update（forward / backward），权重在 trainer 与 rollout worker 间本地同步。该模式适用于小规模、严格 on-policy 实验；在 agentic RL / coding RL 场景中，rollout 的负载形态接近线上 agent serving——请求长、并发高、工具链复杂、长尾明显，需要成熟 serving 的 batching、调度、监控与弹性扩缩。

因此一种架构选择是：将 rollout generation 交由现有生产推理集群或 serverless inference pool 的低优先级可用容量，而非在 trainer 进程内调用 model.generate()。

以 Cursor + Fireworks 为例：Cursor 具备 coding agent 场景与大规模 trajectory 需求；Fireworks 提供生产级推理、弹性 fleet 编排、batching 与低延迟 serving。公开材料显示，Cursor 在与生产 serving 相同的基础设施上运行 RL rollout inference，并可在生产流量低谷时将容量分配给 RL、高峰时收回。

这里的关键不是把 RL 请求当作普通线上请求混入用户路径，而是让 rollout 作为低优先级 / 可抢占 workload 使用生产态推理集群的闲置容量。Fireworks 亦将 rollout inference 产品化：团队自带 trainer，使用 Fireworks 提供 rollout serving、policy publish / version pinning 与 fleet 编排。

这种路线的潜在收益：

• 复用生产 serving 已有的调度、batching、KV cache、监控、健康检查；
• rollout 容量可弹性扩缩，无需与 trainer GPU 绑定；
• 线上 trace、离线 rollout、evaluation traffic 更容易统一在一个 inference API 层；
• agent 应用侧和 RL 训练侧可以共享同一套 LLM proxy / inference gateway 能力，但通过优先级、配额、版本和租户隔离避免相互污染。

代价同样明确：

• policy 权重或版本如何快速、可靠地发布到推理集群；
• old policy / new policy / reference policy 的版本如何隔离与路由；
• logprob、token trace、sampling params 是否能完整回传，满足 on-policy / off-policy correction 需求；
• serverless / production serving 的调度目标往往是 SLA 和租户公平，而 RL 训练目标是 sample goodput，两者可能冲突；
• 网络传输、跨服务 retry 和 inference queue delay 会直接进入 training loop。

不应将生产推理理解为必然优于 colocated rollout，而是 rollout 后端可分为三类：colocated / 同 job 部署（1）、训练专用 rollout 服务（2）、生产推理平台式（3）。1 适用于严格 on-policy 与小规模实验；2 适用于需要独立扩缩但仍专属于某次训练的任务；3 适用于大规模 agentic data generation、evaluation 与近生产流量回流。

调度问题变成 RL 系统的核心能力

服务化之后，调度不再等价于“申请 N 张 GPU”。policy update、rollout generation、reward / judge、environment worker 的资源画像差异很大：

• trainer 需要 GPU 做 forward / backward（及 reference / value 等辅助 forward），对通信拓扑和 checkpoint 敏感；
• rollout 需要 GPU 做推理，对 batching 和 KV cache 敏感；
• reward 可能需要 CPU、GPU 或外部 judge API；
• environment 可能需要 CPU、容器、文件系统、网络隔离。

因此调度层要管理的是 异构 worker 的生命周期：谁启动、谁扩缩、谁与谁做 colocated placement、谁可以异步、谁失败只影响局部。

Ray 调度框架（OpenRLHF / SkyRL / AReaL / ROLL）以 Ray actor / task 组织 rollout、reward、env 等异构 worker。Relax 采用服务化异步与 TransferQueue，支持可配置 staleness。verl 以 HybridFlow 组织 train / generate / reward 节点。slime 谱系（slime / Miles / vime）采用单控制器 + Data Buffer 路径，结构相对直接；复杂 agentic workload 下的 env 编排与异步扩展可能需要额外层。

Colocated 与 Disaggregated 部署是成本和语义的 trade-off

Colocated deployment：训练与 rollout 共用 GPU，权重同步路径短，适用于小集群、同步 RL 与严格 on-policy。代价是 rollout 的 generation forward 与 trainer 的 update forward / backward 竞争 GPU 资源，长尾 generation 易产生 GPU bubble。

Disaggregated deployment：rollout / trainer / reward / env 独立扩缩，适用于大规模与 agentic RL。代价包括 staleness 上升、weight transfer / policy publish 复杂度增加、队列积压与跨服务 failure 需显式处理。Relax 的 TransferQueue、AReaL 的 fully asynchronous、DORA 的 multi-version streaming rollout 均针对 disaggregated 部署下的 goodput 与 off-policy 控制。

若 rollout 接生产推理集群，则为更彻底的 disaggregation：训练系统仅消费 API 返回的 completions、logprobs、trace 与 reward 所需元数据，不再管理 rollout worker 的本地生命周期。未来框架或将 colocated / disaggregated / production-inference rollout / async 设计为 可配置连续谱，而非二选一开关。

调度层需要显式处理 backpressure 和 staleness

图示：队列、Backpressure 与 Staleness

Disaggregated 部署和服务化之后，training loop 的速度由最慢环节决定，而且各环节的积压会改变样本语义：

• rollout 比 update 慢：trainer 空转，GPU 利用率下降；
• rollout 比 update 快：trajectory 堆积，样本变旧，off-policy 程度上升；
• reward / judge 慢：reward 队列阻塞，rollout 完成但无法进入 update；
• environment 阻塞导致整批 rollout 等待，straggler 效应放大。

若 rollout 接生产推理集群，还将叠加 inference queue delay、rate limit、model version rollout delay、跨服务 retry 等问题。系统因此需要显式语义，而非仅扩大队列容量：

• backpressure：下游慢时上游减速或丢弃；
• timeout 与 partial retry：单条 trajectory 失败不阻塞整批；
• priority queue：eval / on-policy 样本与 backlog 样本分级；
• staleness bound 与 policy version tagging：每条样本标记生成时的 policy 版本，update 侧决定是否接受。

异步 RL 因而不仅是 --async 开关，而是 队列语义与版本语义设计。

服务化后的核心指标是 loop goodput

服务化之后，单点峰值吞吐容易误导优化方向。不应只看 trainer tokens/s 或 rollout tokens/s，更应该看 loop 指标：

• valid trajectories / hour；
• accepted reward samples / hour；
• RL step wall-clock；
• p95 / p99 rollout latency；
• timeout rate；
• judge failure rate；
• policy publish / weight transfer latency；
• stale sample ratio；
• inference queue delay；
• model version publish latency；
• recovery time after worker failure；
• teacher feedback latency；
• teacher feedback acceptance rate；
• feedback-token alignment error rate；
• teacher / verifier disagreement rate。

这些指标共同回答一个问题：单位时间内有多少有效训练样本穿过完整 loop。RL infra 的优化目标应是有效训练样本吞吐，而不是某个组件在独立基准测试里的峰值吞吐。

吞吐瓶颈：Rollout 与长尾

观点：后训练 RL 的 goodput 往往被 rollout generation 和环境长尾支配，而不是只被 backward 性能支配。

分析线索：拆解 rollout、reward / judge、policy update、checkpoint、policy publish / weight transfer 的等待关系。

后训练 RL 的成本不仅来自反向传播。对于长上下文、多轮 agent、工具调用任务，rollout 开销常超过 update。

主要原因包括：

• 自回归生成无法完全并行；
• 轨迹长度差异大，长尾请求延长 step 完成时间；
• 工具调用和环境执行会引入外部等待；
• judge / reward 可能需再次调用模型或执行程序；
• policy 变化后，推理引擎需要同步权重；
• 多轮任务会产生更长 KV cache 和更多状态。

因此 RL 框架的核心指标不应仅包含 tokens/s，还需关注：

• 每个 RL step 的 wall-clock；
• rollout 和 update 是否 overlap；
• p50 / p99 generation latency；
• straggler 对 batch 完成时间的影响；
• 失败轨迹是否能局部重试；
• 权重同步或版本发布是否阻塞；
• 环境执行是否可并发扩展。

如果 rollout 是瓶颈，那么优化 trainer kernel 可能收益有限；反而是 speculative decoding、batching、请求切分、环境并发、超时策略、异步 reward、长尾截断更重要。

框架坐标与社区路标

开源与工业路线

本节按技术特征（调度层、训练栈、代码谱系与系统架构）归纳常见开源框架。分类依据是架构与工程取向，不是维护团队或机构；同一路线中的项目可能由不同主体维护。一个框架可能同时满足多种特征（例如 Relax 与 slime / Miles / vime 同属 Megatron + SGLang 取向，但系统架构上更应归为服务化异步路线）。

TRL：Hugging Face 生态下的 post-training 库，支持 SFT、DPO、PPO、GRPO 等。

OpenRLHF：基于 Ray、vLLM、DeepSpeed 的 RLHF 框架，支持 PPO、DPO、KTO、GRPO 等，并提供 single-turn / multi-turn agent 模式。

verl：HybridFlow 实现，将 actor、rollout、reward、trainer 组织为可组合 dataflow，训练后端支持 FSDP / Megatron，rollout 后端支持 vLLM / SGLang。

Ray 调度：以下框架以 Ray actor / task 组织 rollout、reward、env、trainer 等异构 worker。OpenRLHF 侧重 RLHF 工程集成；SkyRL 提供 multi-turn agent RL 与异构调度，SkyRL-Agent 包含 pipeline dispatcher；AReaL 实现 generation 与 training 的 fully async 解耦；ROLL 集成 Ray、Megatron-Core、SGLang / vLLM，支持 Reinforce++、GRPO、GSPO、AutoDeviceMapping。NeMo-RL、Relax 也使用 Ray，但 NeMo-RL 主要归类为生产 infra，Relax 主要归类为服务化异步——「使用 Ray」本身不足以界定框架类别。

slime 代码谱系：以下项目继承 slime 的 Megatron 参数透传、Data Buffer 与低抽象设计，含直接 fork 或基于 slime 训练栈的再实现。slime：Megatron 训练 + SGLang rollout；Miles：slime fork，增加 FP8、R3、speculative RL、VLM multi-turn 等；vime：保留 slime 训练栈，rollout 后端改为 vLLM + vllm-router。Relax 与这一路线同样强调 Megatron + SGLang，但更适合按 Ray Serve + TransferQueue 的服务化异步路线理解。

DeepSpeed-Chat：较早的 DeepSpeed-based RLHF pipeline，将 SFT、reward model、RLHF / PPO 等阶段串成示例化流程。其定位接近早期工程参考，不宜与 modern vLLM / SGLang rollout service 或 Ray 异构调度框架等同视之。

NeMo-RL：NVIDIA 生产向 RL 框架，支持 Ray、Megatron、vLLM / SGLang、checkpoint、拓扑放置、异步 rollout、environment isolation、FP8 等。

Forge：MiniMax 内部 agent-native RL 框架（未开源）。公开材料描述三层结构：Agent Side、Middleware（Gateway + Data Pool）、Training-Inference Side（Rollout Engine + Train Engine）。支持黑盒 Agent 通过推理 gateway 接入；算法包括 CISPO、process reward、task completion time reward；工程特性包括 windowed-FIFO 调度、prefix tree merging、MTP speculative decoding、global KV cache pool。

补充分类

上一节按技术谱系列举主要 repo。以下按用途与架构模式补充几类，并说明部分共性。

工业向框架的共性：verl、OpenRLHF、slime、ROLL、NeMo-RL 等均在演进 rollout 独立部署、reward 模块化、MoE / FP8 / async 支持。差异在于调度层选型（Ray vs dataflow vs 单控制器）和训练 / rollout 后端组合。

研究向框架：Safe-RLHF 面向 helpfulness / harmlessness 约束下的 RL 对齐；trlX 曾提供 distributed RL for LLM 能力，近年维护活跃度较低。

Agent 接入的两种模式：

Agent Lightning 支持 RL、APO、SFT 等，LightningRL 将 agent 执行建模为 MDP 并做 credit assignment。rLLM 通过 @rllm.rollout 或 LiteLLM proxy 从 vLLM 获取 token id 与 logprob，训练后端为 verl 或 tinker；v0.2+ 提供 AgentWorkflowEngine，支持多 agent workflow。两种零侵入方案均需处理 agent 侧与训练侧 tokenizer / chat template 不一致带来的 logprob 与 mask 对齐问题。

轻量 / 实验向：Unsloth 提供单卡 GRPO / DPO 优化；LLaMA-Factory 支持 100+ 模型的 SFT / DPO / PPO；EasyR1 为 verl fork，面向 Qwen-VL 等多模态 GRPO、DAPO、LoRA RL。通常用于算法与 reward 验证，再迁移到 verl / slime 等大规模栈。

代码 fork 与衍生：slime 谱系的 fork（Miles、vime）共享 Megatron 参数透传与 Data Buffer，差异在 rollout 后端（SGLang / vLLM）与工程侧重（FP8 / R3、vLLM 生态）。Relax 与该谱系在 Megatron + SGLang 取向上相近，但核心差异是 Ray Serve + TransferQueue 带来的服务化异步架构。verl 谱系的 fork（EasyR1、VeRL-Omni、uni-agent）共享 HybridFlow，分别面向多模态 RL、diffusion / omni-modal、agent env 统一层。

框架可按以下维度区分定位，无单一排序标准：

选型时可按 workload 与 infra 条件匹配，而非按框架名称：

• preference optimization 实验：TRL、Unsloth、LLaMA-Factory
• 多模态 RLVR：EasyR1、VeRL-Omni
• 可扩展 RLHF：需关注生成后端、训练后端、权重同步
• 已有 vLLM infra：vime、OpenRLHF、rLLM；已有 SGLang + Megatron：slime、Miles、Relax
• agentic RL，scaffold 固定且可控：Relax、OpenRLHF、uni-agent 的内建 env
• agentic RL，scaffold 异构或黑盒：Agent Lightning、rLLM、Forge gateway（公开材料）
• 大规模 MoE：需关注拓扑、checkpoint、failure recovery
• 大规模异构 Agent scaffold：需关注 Middleware 解耦、trajectory 协议、async 调度

社区路标

观点：RL 框架分析不应仅关注当前功能列表，还需结合社区路标判断下一阶段抽象的可能收敛方向。

分析线索：持续跟踪 README news、release notes、experimental 模块、issues、meetup / tech blogs 与论文代码。

分析 RL 框架时，还需系统收集社区路标。许多项目无正式 ROADMAP.md，方向通常分散于 README news、release notes、experimental 模块、issues / discussions、meetup / tech blogs 与论文代码。

从当前社区信号看，以下主线较为清晰。

单轮 RLVR 到多轮 Agentic RL

verl 已将 multi-turn with tool calling、search tool、sandbox integration 纳入文档，experimental 方向包括 uni-agent、VeRL-Omni、fully async policy、one-step off-policy、VLA 等。

OpenRLHF README 将 single-turn agent 与 multi-turn agent 分为两种执行模式，并强调 RL algorithm 与 agent executor 解耦；multi-turn 路线包括 custom agent function、环境反馈、async pipeline 与 VLM multi-turn RL。

Forge 进一步扩展 agent executor 解耦：Agent Side 与 Train/Rollout Side 通过 Gateway + Data Pool 隔离，白盒 Agent 可将 Context Management 建模为环境状态转移，黑盒 Agent 将请求路由至 RL Gateway，由框架在推理侧采集 trajectory。MiniMax 公开材料称 M2.5 训练接入了 10 万+ 种 scaffold，multi-turn agent RL 的瓶颈正从 tool 接入能力转向异构 scaffold 的规模化接入。

Agent Lightning 与 rLLM 表明：在 episode 抽取可靠的前提下，无需预先修改 agent 代码。Agent Lightning 支持 RL、APO、SFT 等算法，Youtu-Agent 在其修改分支上验证了 128 GPU 稳定收敛；rLLM 提供 @rllm.rollout + verl 后端的 CLI-first workflow，内置 50+ benchmark。社区方向正从框架内 multi-turn env 扩展至更广泛的 agent 可训练性。

社区正在将 RLVR 从「prompt → answer → verifier」扩展为：

1

observation -> action -> environment feedback -> next action -> final reward

算法名称仍可沿用 GRPO / PPO / REINFORCE++，但样本形态已发生变化。

Rollout 后端服务化

Rollout 不再只是 trainer 里的一段 generate()，而是一个可以独立优化、扩容、观测和复用的系统。

OpenRLHF 已支持 async training、partial rollout、vLLM pause/resume、OpenAI-compatible agent server executor；verl 强调 vLLM / SGLang / HF Transformers 多后端及 train / generate 间的 resharding。

社区论文中亦出现 ProRL Agent 等 Rollout-as-a-Service 设计。rollout 生命周期从 trainer 中拆分 可独立优化与复用；否则多环境、多工具、多沙箱、多模型版本将使训练代码复杂度持续上升。

异步 RL：性能方向，但存在 trade-off

多个框架都在走 async：

• verl 的 experimental 里保留 fully async policy、one-step off-policy；
• OpenRLHF 支持 async RLHF / async agent RLHF，也明确提示 async queue 会带来 off-policy 程度；
• NeMo-RL 的功能列表里写了 Async RL、asynchronous rollouts、replay buffers、fully asynchronous GRPO；
• Forge 用 windowed-FIFO 在吞吐与分布一致性之间折中：窗口内允许局部乱序以消除 straggler，窗口外禁止跳读，避免训练分布被 fast sample 偏移。

为提高 goodput，社区持续将 rollout、reward、training 进行 overlap。

但异步亦引入新问题：样本可能来自旧 policy，in-flight rollout 可能混合新旧权重，advantage / KL / correction 的设计因而更为关键。async RL 因此不仅是 --async 开关，还需配套 off-policy correction、replay buffer、staleness control、partial rollout consistency 等机制。

多模态与工具环境进入主线

TRL 定位为 Hugging Face 生态的 post-training trainer 集合，覆盖 SFT、GRPO、DPO、RewardTrainer 等入口。

系统型框架正在扩展多模态与环境集成：

• verl 新闻里有 VeRL-Omni、VLM / multi-modal RL；
• OpenRLHF 0.10 增加 VLM RLHF 和 Multi-Turn VLM RL；
• NeMo-RL 支持 VLM、SGLang rollout、NeMo-Gym、multi-turn RL、environment isolation。

这意味着后训练 RL 的“环境”不再只返回文本。后续可能会出现图像 observation、浏览器截图、视频帧、GUI state、代码执行结果、数据库查询结果、文件 diff 等混合状态。

大规模生产化压过单点算法

NeMo-RL 的功能列表 / roadmap 侧重 infra：SGLang backend、speculative decoding、Yarn long-context、Megatron inference、fault tolerance / auto-scaling、MoE performance、FP8、GB200、environment / worker isolation。

Forge 的 roadmap 侧重 agent scaffold 规模化：黑盒 / 白盒 Agent 统一接入、Context Management 纳入 RL 交互环、prefix tree merging、heterogeneous PD disaggregation、global L3 KV cache pool，以及 mixed-domain（Reasoning / General QA / Agent）联合训练。工业实践中，scaffold 支持范围与内置训练目标种类同为框架能力的重要维度。

verl 也在 FSDP2、Megatron backend、LoRA RL、large MoE、AMD ROCm、SGLang unique features 上持续推进。

大规模 RL 框架的重点将继续包括：

• rollout 加速；
• 训练 / 推理权重转换或同步；
• MoE 与长上下文性能；
• 多硬件支持；
• checkpoint / fault tolerance；
• 环境和 worker 隔离；
• recipe 可复现。

社区路标可区分为算法路标与系统路标：前者关注优化方法，后者关注能否持续、低成本、稳定地完成训练。

社区原型先于主框架暴露需求

除主框架外，还需关注其生态中的衍生项目。

VerlTool 等工作在 Verl 生态中集成 tool use、异步 rollout、多模态 observation 与统一工具 API，反映 tool-augmented RL 对插件化与环境标准化的需求。

AgentRL 强调 multi-turn、multi-task、fully asynchronous generation-training pipeline、containerized environment 与 centralized controller，反映 agentic RL 在 multi-task 环境与异步流水线方面的工程压力。

MUA-RL 将 simulated user 纳入 RL loop，表明 multi-turn agent 除 tool calling 外，亦涉及动态用户意图与澄清过程。

这些项目未必都会成为最终标准，但它们能提前暴露主框架即将补的接口：tool API、environment controller、simulated user、sandbox lifecycle、multi-task normalization、cross-policy sampling、rollout service。

因此跟踪 RL 框架时，不应仅关注 release tag，可维护如下社区路标表：

补充展开：Objective、Agentic RL 与系统扩展

前文已从系统视角、Runtime 组件和框架坐标说明：后训练 RL 的难点不只在 objective。本章补充三条更细的逻辑链：

• objective / feedback 路线如何从 PPO、GRPO 扩展到 OPD；
• agentic RL 为什么会放大 credit assignment、token 对齐和稳定性问题；
• 这些问题如何反映到同步 / 异步、colocated / disaggregated 部署和大规模 goodput 上。

Objective 演进：从 PPO 到无 critic RLVR

PPO 的价值是给 RLHF 一个稳定的“多轮复用 rollout”机制：clip ratio、value model、entropy bonus 共同控制单步更新幅度，适合偏好模型打分的 chat alignment；但在 RLVR reasoning 场景，critic 成本和不稳定性会被放大。DPO / KTO 采用离线 preference optimization 路径，省去在线 rollout，适用于快速实验，但不直接覆盖需环境交互的 agentic RL。

GRPO 的关键转向是去掉 value model，用同一 prompt 下多条采样的组内相对 reward 做 baseline，DeepSeek-R1 证明了它可以支撑 rule-based reward 的大规模 reasoning RL。RLOO（Cohere, 2024）和 ReMax（ICML 2024）也是无 critic 路线：前者用 leave-one-out baseline，后者用 greedy baseline；和 GRPO 接近，但公开讨论较少。REINFORCE++（OpenRLHF, 2025）强调全局 batch 归一化，将问题表述为低耦合 policy gradient，但方差与样本效率压力仍存在。

DAPO 的动机是复现级工程：在 DAPO 中，decoupled clip、dynamic sampling、token-level policy gradient、overlong reward shaping 等机制旨在提高进入梯度的有效样本比例，其 Qwen2.5-32B 在 AIME 2024 达到 50 分，表明算法细节与数据过滤同等重要。

CISPO 采用不同路径：在 MiniMax-M1 中放弃 trust region 约束，改为 clip importance sampling weights，让所有 token 都参与梯度；Forge 把它继续用于长 horizon agent RL，并叠加 process reward 与 completion time reward。GSPO 再进一步把优化粒度从 token ratio 推到 sequence likelihood，GSPO 声称 sequence-level clipping 能提升效率并稳定 MoE RL。

DrGRPO 则是对 GRPO 长度偏置的修正，Understanding R1-Zero-Like Training 指出 GRPO 会鼓励错误长答案变长，并用 DrGRPO 改善 token efficiency。脉络上看，算法不是越来越复杂，而是围绕“方差、长度偏置、有效梯度、MoE 稳定性、系统可实现性”反复改写优化粒度。

OPD（On-policy / Online Policy Distillation）：RLVR 与 distillation 之间的第三条路线

若将 post-training 划分为三条数据流，差异在于 feedback 来源与密度，而非是否存在 rollout：

• RLVR / PPO / GRPO：学生 rollout → verifier / reward → policy gradient（A · log π 及 clip、KL）
• SFT / offline distillation：教师数据或教师输出 → 离线监督，但 teacher 分布与 student 当前分布不一致，存在 off-policy distribution shift
• OPD（On-policy / Online Policy Distillation）：学生 rollout 自己的轨迹 → teacher 在这些 on-policy 轨迹上提供密集反馈 → distillation update

OPD 的动机来自两条路线的各自短板。SFT 和 offline distillation 的问题是：teacher 数据来自 teacher 自己的分布，学生一旦偏离训练分布，纠错信号就变弱——distribution shift 会削弱纠错信号的有效性。PPO / GRPO / RLVR 则相反：on-policy 数据生成可行，但 reward 常稀疏、方差高，长链推理与 agent trajectory 中 credit assignment 困难——模型仅知最终失败，难以定位需修正的中间步骤。OPD 结合两者：由 student 执行 rollout，在 student 实际到达的状态上由 teacher 提供密集监督。

与 RL 的关系需单独说明。OPD 保留 on-policy 数据生成，可缓解 off-policy distribution shift；但通常不直接优化 A · log π，也不一定采用 value model、advantage、KL penalty 等 RLHF 结构。其形式更接近将 teacher feedback 作为 dense reward / dense target，以 teacher 局部偏好替代稀疏 verifier 信号，降低 credit assignment 难度。因此 OPD 不宜简单归为 PPO / GRPO 变体，而应视为连接 RL、distillation、reasoning supervision 的 post-training 路线。

2026 年已有若干工作从不同侧面验证这条路线。DiffusionOPD 将 OPD 扩展至 diffusion model 的多任务 RL / 后训练：task-specific teachers 先独立学习，再沿 student rollout distill 至统一 student，表明 OPD 不局限于离散 token LLM，亦可推广至连续状态生成模型。

OmniOPD 提出 logit-free OPD，用 chunk-level semantic verification 替代直接 teacher logits——teacher 无需暴露 token logits，黑盒 teacher 或语义验证亦可构造密集反馈。SG-OPD 用 verifier 作为 teacher feedback 的 trust signal，表明 token-level preference 并非总可靠，需 gating 或 verifier 校准。

POPD / TOPD（Are Full Rollouts Necessary for On-Policy Distillation?）指出 OPD 不一定需要完整 rollout：Progressive OPD 和 Truncated OPD 可控制 rollout horizon，论文报告 TOPD 只用 10% rollout horizon 可接近 OPD 效果，POPD 最高带来 3x 训练效率提升。CaMOPD 做 Multi-Teacher On-Policy Distillation，用于 general capability recovery 与 domain preservation，表明 OPD 可承载多 teacher 能力整合，但需处理不同 teacher / domain 的梯度冲突。

OPD 的关键变量包括 teacher feedback 的形态——token-level logits、chunk-level semantic feedback、verifier-gated feedback、multi-teacher feedback、truncated / progressive horizon feedback。形态不同，系统接口、延迟预算与稳定性策略亦不同。

亦需说明：OPD 并非对所有场景均适用，它会把 RLVR 中的 reward bottleneck 转移为 teacher service bottleneck——teacher 推理成本高；teacher logits / chunk feedback 难以低延迟回传；teacher 与 student 的 tokenizer / chat template 不一致会带来对齐问题；teacher feedback 可能在学生错误轨迹后半段失真；多 teacher 可能产生梯度冲突。工程上，OPD 的收益取决于 teacher 服务是否具备与 rollout 同等的扩展、观测与版本化能力。

图示：OPD Teacher / Distillation Service

Agentic RL 的训练语义挑战

Agentic RL 的首要难点是多轮 credit assignment：最终 reward 可能仅在任务完成时出现，中间则包含 search、tool call、file edit、terminal execution、browser observation 等步骤。将整条 trajectory 拼接为单一 sequence 会混淆语义边界；tool observation 不宜作为 model action 参与 loss 计算。

框架内路线与零侵入路线正在形成对照组。Relax 在工程层提供 BaseInteractionEnv、multi-turn loss masking、VLM context carry-over，将 execute → observe → decide 抽象为框架原语。SkyRL-Agent 表明 agent rollout 需专用 dispatcher：SkyRL-Agent 在 SWE-Bench Verified 上将 Qwen3-32B 从 24.4% Pass@1 训练至 39.4%，并报告异步 pipeline dispatcher 比 naive async batching 快 1.55x。

零侵入路线将问题表述为 如何从任意 agent 代码中稳定抽取 training transition。Agent Lightning 的 LightningRL 将 agent 执行建模为 MDP，引入 credit assignment module，将 agent 轨迹拆分为 training transition；其价值在于 agent execution 与 training 的 disaggregation，LangChain、AutoGen、OpenAI Agents SDK 等可通过 minimal-change 接入。rLLM 采用相同思路，并强调 inference-native token capture：LiteLLM proxy 或 SDK 从 vLLM 获取 token id 与 logprob，写入 SQLite store，再交由 verl 后端更新——同一 agent 代码可用于 eval 与 train。rLLM 的 @rllm.rollout + AgentWorkflowEngine 表明零侵入方案亦适用于 solver–judge、planner–executor 等多 agent workflow。

Retokenization drift 是零侵入接入中的常见问题：agent 侧 API tokenization、训练侧 tokenizer、工具模板和 chat template 不一致时，logprob / mask / reward 对齐会发生漂移。rLLM 的应对是从 inference server 侧取 token id，Agent Lightning 强调 episode schema 与 credit assignment，框架内路线靠 loss masking 和 env observation 边界。三条路线最终都要收敛到同一数据结构：显式区分 model action 与 environment observation，并保留原始 message 与 token span。

稳定性挑战：信号、策略与轨迹的失配

entropy collapse 指 policy 过早确定化：reward 稀疏而更新置信度高时，分布熵持续下降，探索减少、梯度减弱，policy 收敛至局部策略。OPEFO 从 token-level entropy flow 解释该现象，认为 entropy-decreasing tokens 长期主导 entropy-increasing tokens，需按熵贡献重标定更新；另有工作通过 dynamic clipping 调控 entropy。

reward hacking 属于 Goodhart 问题：verifier、unit test、LLM-as-judge 均为代理目标，模型可能优化代理指标而非任务本身；在代码与 agent 任务中表现为修改测试、利用答案泄漏、生成格式正确但语义无效的输出。KL 异常通常来自 reward 过强、reference 约束过弱或采样分布突变；KL 仅约束分布距离，不约束工具权限与任务语义。

长尾生成兼具系统稳定性与算法层面含义：长输出可能包含有效推理，也可能来自过度思考或长度投机；FP8-RL 将 rollout 视为瓶颈，W8A8 + FP8 KV-cache + TIS/MIS correction 报告最高 44% rollout throughput gain。大 MoE RL collapse 还涉及 expert routing、低精度 rollout、train-inference mismatch 与负载不均；GSPO、FP8 correction、R3/Routing Replay、Relax 的 R3 支持均旨在控制 routing 与 rollout 偏差对 policy update 的放大效应。

OPD 亦可从稳定性角度分析：其缓解 RLVR 的 sparse reward 与高方差，但引入 teacher-feedback stability 问题。标准 OPD 假设 teacher 在学生轨迹上的 token-level preference 可靠；学生早期轨迹质量通常较低时，teacher 对后续 token 的局部 logits 未必指向有效修正方向——feedback 可能含噪或产生误导。SG-OPD、OmniOPD、TOPD/POPD 分别从 verifier gating、logit-free chunk feedback、rollout horizon control 三个方向处理该问题，表明 OPD 工程化需显式设计 feedback 可信度与 horizon 边界，而非仅替换 teacher。

系统架构权衡：同步、异步与部署边界

同步训练的优点是 on-policy 语义明确：一批 rollout 来自同一 policy，KL、advantage、clip 的解释较为直接；slime / vime 等单控制器路线将 Megatron 训练、SGLang 或 vLLM rollout、Data Buffer 置于显式路径，调试与正确性验证相对清晰。缺点是长尾 generation 导致整批等待，rollout 越长、工具越慢，GPU bubble 越显著。

异步路线的代表为 AReaL 与 Relax：AReaL 完全解耦 generation 与 training，报告最高 2.57x speedup；Relax 通过 TransferQueue 与 configurable staleness，覆盖 on-policy 至 fully async 的连续谱，论文报告 Qwen3-Omni-30B fully async 相比 colocated deployment 2.00x speedup。

colocated deployment 适用于小集群与严格 on-policy，actor / rollout 复用 GPU，权重同步简单；disaggregated deployment 适用于大规模与 agentic workload，rollout、reward、env 可独立扩缩，但需处理 staleness、policy publish / weight transfer 与故障隔离。SGLang 侧重结构化、多轮、KV 复用与复杂程序执行，SGLang 报告最高 6.4x throughput；vLLM 侧重通用 serving、PagedAttention 与生态覆盖，vLLM 报告 2–4x throughput。框架选型应基于 workload 是单轮大批量、长上下文多轮还是工具环境驱动，而非单一吞吐 benchmark。

扩展规律：从单点吞吐到 loop goodput

后训练 scaling 与预训练不同：预训练主要是 token、参数、算力的平滑缩放；RL post-training 是 policy、rollout、reward、environment、judge、checkpoint、policy publish / weight transfer 的全链路缩放。扩展瓶颈也不是单个 GPU TFLOPS，而是 loop goodput：多少有效 trajectory 能在单位时间内产生、被打分、被正确归因并进入更新。

OpenRLHF 早期就指出 RLHF 不是单模型训练，而是四模型协调问题，因而用 Ray、vLLM、DeepSpeed 重做 70B+ 调度。HybridFlow / verl 将 RLHF dataflow 中各 NN 节点扩展为分布式训练或生成程序，并以 hybrid controller 表达依赖，报告 1.53x–20.57x throughput improvement；核心是 3D-HybridEngine 降低 train / generate resharding 成本。

NeMo-RL 的功能列表 / roadmap 面向千卡生产：NeMo-RL 支持 DTensor、Megatron、vLLM、SGLang、Megatron inference、FP8、Async RL、environment isolation，NeMo-Aligner 时代已披露千 GPU 规模下 340B / 405B 级模型训练能力。AReaL-Hex 指出异构 GPU 下 rollout generation、reward、policy update 的 compute / memory / communication profile 不同，调度需按 stage 匹配资源，论文报告同预算最高 1.50x throughput、等吞吐最高 1.46x cost reduction。

概括而言，RL scaling 更接近全链路吞吐规律，受最慢环节、长尾、staleness 与失败恢复共同约束。

生态格局：工程路线分化，核心抽象收敛

上文「开源与工业路线」已按技术特征梳理各框架。本节不再重复框架清单，而是从演化路径看：当前生态更像四条路线并行，而不是单一框架收敛。

slime 谱系的技术哲学是少抽象、强原生：slime 直接暴露 Megatron 与 SGLang 参数，强调 Data Buffer、workflow、reward / verifier 的显式路径；Miles 沿此路线补 FP8、R3、speculative RL 与企业级 MoE 稳定性；vime 则把 slime 训练栈横向接入 vLLM + vllm-router 生态。这一路线适合已有 Megatron 训练能力、需要低层可控性的团队，代价是 agent environment、异步调度和服务治理通常要在外层补。

verl 谱系更像可组合 dataflow 承载层：HybridFlow / 3D-HybridEngine 让 train、generate、reward 节点可以按依赖组合，EasyR1、VeRL-Omni、uni-agent 分别把它推向多模态 RLVR、diffusion / omni-modal、agent env 统一层。它的优势是研究原型和系统扩展都能挂在同一抽象下，代价是抽象层较重，深度定制时需理解 controller、worker、后端 resharding 的边界。

Ray 系（OpenRLHF、SkyRL、AReaL、ROLL）强调异构 worker 调度与执行弹性，适合 rollout、reward、environment、trainer 分工复杂的 workload。Ray actor / task 让多角色生命周期更自然，但正确性不会自动获得；policy version、staleness、queue backlog、partial retry、跨服务观测仍需框架显式建模。

独立工业 / infra 路线继续分化：Relax 代表 Ray Serve + TransferQueue 的服务化异步；NeMo-RL 绑定 NVIDIA 硬件、Megatron、DTensor、checkpoint 与 worker isolation；ROLL 是 Ray + Megatron-Core + vLLM / SGLang 的全栈路线；Forge 公开材料体现 agent-native gateway、Data Pool 与 scaffold 解耦。

这些路线未必会收敛到相同代码库，但会在 episode schema、reward contract、rollout service、policy publish / weight sync、environment isolation 与 observability 上趋同。真正继续分化的是训练后端、推理后端、调度控制面和组织已有 infra 的方式。

趋势收束：下一步值得关注什么

从以上链条看，未来 12 个月可能广泛采用的不是单一方法，而是四类能力：异步 agent RL、过程奖励模型（PRM）、OPD / on-policy distillation，以及生产态推理集群上的低优先级 rollout。PRM 针对 reward 稀疏，异步 agent RL 针对 rollout goodput，OPD 通过 teacher dense feedback 处理长程 credit assignment 与 RL 方差，生产态 rollout 则针对大规模生成的资源弹性与成本。

OPD 可与 PRM、verifier、black-box teacher 组合：PRM 提供过程性判断，teacher 提供局部改写方向，verifier 提供 gating。其定位可能介于 RLVR 与 distillation 之间。异步 agent RL 已有 AReaL、Relax、SkyRL-Agent、DORA 等系统实践：DORA 报告 rollout 占 step time 50–80%，multi-version streaming rollout 在保持 intra-trajectory policy consistency、data integrity、bounded staleness 的同时消除 bubble，系统吞吐提升 2–3x、工业规模加速 2–4x；此类方案要求框架原生支持 policy version 与 staleness 语义。

PRM 可作为 reward hacking 的补充：纯 outcome reward 稀疏，LLM-as-judge 存在 exploit 风险，过程奖励可将长链推理、工具调用、代码修改分解为更细粒度监督，代价包括标注、judge 稳定性与过程过拟合。生产态 inference rollout 可能先在 coding agent、长程 evaluation、synthetic data generation 中落地，因为这些 workload 对 p99 用户延迟不敏感，但需要大规模并发生成；关键约束是 priority isolation、policy version pinning、logprob / trace export 与 preemption 后的样本完整性。

多智能体 RL 可能在 communication cost、self-play、simulated user 等方向先出现，Agent-GSPO、MUA-RL 表明 token economy 与动态用户意图将进入 reward 设计。omni-modal RL 持续扩展，Relax 与 NeMo-RL 已将 VLM / omni-modal 纳入主线；VeRL-Omni 表明 diffusion / omni-modal post-training 可与 LLM RL 共享部分 orchestration。

零侵入 Agent RL（Agent Lightning、rLLM）与框架内 agent 层（uni-agent、Forge gateway）并行发展：前者适用于 scaffold 规模大、需复用现有 agent 资产的场景；后者适用于 env 可控、内循环吞吐要求高的场景。diffusion model RL 可能在图像 / 视频生成质量与偏好优化中先行落地，与 agentic RL 的统一仍需 episode 与 reward 抽象进一步泛化。

持续跟踪时，可将重点放在这些项目是否补齐共同抽象：DCPO、DORA、SkyRL-Agent、NeMo-Aligner / NeMo-RL、Forge / CISPO、FP8-RL、EasyR1、VeRL-Omni、Miles R3、vime train–inference parity、production-inference rollout。

落地：目标架构与选型

目标架构：四层分工

观点：未来后训练 RL 系统不会只有一个 trainer 层，而会自然分成应用、环境 / reward、RL orchestration、model runtime 四层。

分析线索：从四层架构看每层职责和边界。

后训练 RL 系统可划分为四层：

1
2
3
4
5
6
7
8
9
10
11

Application / Agent Layer
  任务、工具、环境状态（Env State：文件系统 / 沙箱 / 浏览器等）、用户交互、线上 trace

Environment / Reward Layer
  sandbox、verifier、judge、reward contract、eval harness

RL Orchestration Layer
  rollout、trainer、reference、reward、数据流、checkpoint、调度

Model Runtime Layer
  distributed training、inference engine、KV cache、kernel、通信、拓扑

分层后，各层边界更为明确。

Application 层中的 Env State（环境状态，有时也称 workspace）是 agent 执行任务时可读写的状态容器：coding 场景下是代码仓库与挂载目录，terminal 场景下是沙箱 cwd，browser 场景下是页面 DOM 与交互状态。它与 Task（目标）、Tool（可调用能力）并列，是 tool 作用对象与中间 observation 的载体，不在主训练数据流链上，但为 Rollout 提供执行上下文。

第一，Agent 应用不宜直接绑定特定 trainer，而应以标准 trace 记录任务轨迹、工具调用、观察结果与最终 artifact，供 RL 层消费。

第二，reward / verifier 应作为独立资产，具备版本化、可回放、可审计的 domain contract 属性，而非嵌入单次训练脚本的条件逻辑。

第三，RL orchestration 不应假设 rollout 仅为本地 model.generate()，而需处理 policy serving 更新、异步请求、权重同步 / 版本发布、失败重试与环境并发。

第四，model runtime 层持续专门化，并同时覆盖训练与推理两侧：Training Engine（Megatron / FSDP 等，供 Trainer 执行 update forward / backward）与 Inference Engine（vLLM / SGLang 等，供 Rollout 执行 generation forward）；MoE、长上下文、MTP / speculative decoding、FP8 / FP4、context parallel、expert parallel、KV cache 与通信拓扑均影响 RL goodput。

框架选型：从瓶颈出发

观点：框架选型应该从当前瓶颈出发，而不是从框架名或算法名出发。

分析线索：用 checklist 判断当前系统缺的是 rollout、environment、reward、observability 还是大规模 runtime。

评估 RL 框架或内部系统时，宜优先检查以下问题：

1. Rollout 怎么做
   是否支持 vLLM / SGLang 等高吞吐后端？权重同步 / 版本发布是阻塞还是异步？长尾 generation 怎么处理？

2. 环境怎么接
   是否能表达多轮状态、工具调用、文件系统、终端、浏览器、异常和 timeout？环境失败是否会导致整轮训练中断？

3. Reward 怎么版本化
   reward breakdown 是否记录？judge 版本是否入库？同一批 trajectory 是否可用新 reward replay？

4. 数据结构是否足够完整
   是否保存 old logprob、ref logprob、采样参数、policy checkpoint、trace id、env version？是否可复现指定 trajectory？

5. 异步边界在哪里
   rollout 与 update 是否可 overlap？reward / judge 是否可异步？checkpoint 是否阻塞主循环？

6. 资源放置是否可控
   训练、推理、reward、环境 worker 能否分别申请资源？MoE / EP / TP / PP 是否能感知拓扑？

7. 失败是否局部化
   单条 trajectory 失败、单个 sandbox 阻塞、单个 vLLM worker 异常时，是局部重试还是整个 job 重启？

8. 观测是否面向 RL 训练链路
   除了 loss、reward 均值，有没有 p99 latency、timeout rate、tool error taxonomy、reward distribution drift、policy version 对比？

9. 算法改动成本多高
   改 advantage、KL、reward shaping、loss aggregation、采样策略时，是改一个模块，还是要穿透整个框架？

10. 上线回流路径是否存在
    线上 agent trace 是否可进入离线训练？训练中 eval bad case 是否可回流至数据构造？

11. Rollout / reward / env 是否能独立扩缩
    rollout generation、reward / judge、environment worker 能否分别申请资源、独立扩容和故障隔离？还是必须随 trainer job 同生共死？

12. Rollout 是否必须绑定训练集群
    rollout 是否只能接 job 内嵌的 vLLM / SGLang worker，还是可以接生产态 inference gateway 或 serverless inference pool？若接外部推理平台，policy publish / version pinning 和 trace 回传路径是否明确？

13. 是否记录 policy version、sample staleness、queue delay
    每条 trajectory 是否带 policy version tag？样本进入 update 时 staleness 是否可度量？rollout / reward / inference 各段 queue delay 是否入库？

14. 是否支持 backpressure、partial retry、timeout cleanup
    下游 reward 或 judge 慢时，rollout 是否会 backpressure？单条 trajectory 超时能否 partial retry 或局部丢弃，而不是阻塞整批？

15. 调度层是否能区分异构 worker
    trainer GPU、rollout GPU、CPU env worker、judge worker 能否分别放置、扩缩和监控？placement 是否感知拓扑和资源画像差异？

16. 复用生产推理集群时，训练所需字段能否完整回传
    若 rollout 走外部 inference API，logprob、sampling params、token trace、model version 是否能稳定返回，满足 advantage 估计和 off-policy correction？

17. 是否支持 Teacher / Distillation Service
    OPD 场景下，框架是否能在 student rollout 轨迹上调用 teacher，并回传 token / chunk / step 级反馈？是否支持 teacher version、feedback span、teacher-student tokenizer mapping？

18. OPD feedback 是否可观测和可回放
    teacher feedback、verifier gating、teacher / verifier disagreement、feedback latency 是否入库？同一批 trajectory 能否用不同 teacher 或不同 feedback policy replay？

结论

观点：后训练 RL 的长期积累将沉淀于系统资产，而非单个 trainer repo。

分析线索：从链路质量、runtime 化、agentic 生产系统化、可验证 reward、goodput 五个角度收束。

第一，后训练 RL 的壁垒正从单个算法转向链路质量。算法公开速度快，但高质量环境、reward contract、trace schema、failure taxonomy 与 replay / eval 体系复制成本较高。

第二，RL 框架的竞争维度正从内置训练目标扩展至训练链路组织能力。PPO / GRPO 为入口，难点在于 rollout、reward、environment、trainer、serving、checkpoint、observability 在大规模下的协同稳定性。

第三，agentic RL 推动训练系统接近生产系统。环境、工具与失败模式均为真实场景，除 GPU 计算外，还需处理权限、隔离、超时、状态污染、外部依赖与复现性。

第四，可验证 reward 将成为 post-training 的核心资产。其效果受 verifier 设计约束：verifier 越强，模型越可能围绕 verifier 优化；verifier、hidden eval、过程审计、人工抽检与 reward 多样性需一并设计。

第五，RL infra 的优化目标应为 goodput，而非峰值吞吐。大规模 RL job 的主要成本来自等待、重启、长尾、同步、I/O、checkpoint 与环境失败；降低失败率、缩短恢复时间、支持归因通常比单点 benchmark 更具实际价值。

概括而言：

LLM post-training RL 的成熟实践通常不只维护 trainer repo，而维护任务环境、reward contract、trajectory store、RL runtime 与 eval / replay 组成的系统。

RL 框架仅为其中一层；主要积累在完整链路中。

附录

后训练 Objective 速查

社区和框架文档常用方法名（DPO、GRPO）指代一套 post-training 配方。落到 trainer，差别在 objective 怎么写、要不要在线 rollout、feedback 来自 reward 还是 teacher，以及 advantage / KL / clip 如何进梯度。粗略可分三类：

• Preference objective（DPO、KTO）：直接在 preference 或 good/bad 信号上优化 logprob，不跑 rollout
• Policy-gradient objective（PPO、GRPO 等）：先采样 trajectory，再按 reward 估 advantage，以 A · log π（及 clip、KL）做 backward；PPO 另训 value head，GRPO 及变体去掉 critic
• On-policy distillation objective（OPD）：学生按当前 policy rollout，teacher 沿学生轨迹提供 token / chunk / step 级 feedback，优化形式更接近 dense distillation target

下表按方法名汇总各 objective / feedback 路线的关键形式。SimPO、ORPO、RAFT 等变体未列，差异多在 advantage baseline、KL 注入位置、clip 粒度与 feedback 密度。

OPD 与 RLVR：二者均使用 student 的 on-policy trajectory。RLVR 主要从 reward / verifier 获得稀疏或标量反馈；OPD 从 teacher 获得更密集的 distillation 信号，形式为 on-policy 数据生成 + distillation objective。

这篇是一次代码阅读笔记，起点是一个很具体的问题：

NeMo-RL 怎么将 rank placement 对齐 NVLink domain？

一开始容易把几件事混在一起：

1. Kubernetes/KAI 把 pod 放到同一个 GPU clique / ComputeDomain。
2. Ray worker 启动时注册 nvlink_domain_<UUID> 和 topo_rank 这类 custom resource。
3. NeMo-RL 的 RayVirtualCluster 创建 placement group，并把 Ray bundle 映射成 PyTorch/Megatron rank。
4. Megatron/DeepEP 的 HybridEP 还会读 NUM_OF_HYBRID_EP_RANKS_PER_NVLINK_DOMAIN、USE_MNNVL 这类环境变量。

它们都和“拓扑”有关，但不是同一层的东西。尤其是：当前 main 分支和一些特性分支的行为并不一样。

结论

当前 main 分支里：

• infra YAML 会在 Ray worker 启动时注册 nvlink_domain_<UUID> 和 topo_rank。
• 但是 NeMo-RL Python 侧的 rank placement 没有消费 nvlink_domain_*。
• unified placement group 场景下，RayVirtualCluster 只按 (node_id, gpu_id) 对 bundle 做稳定排序。
• 真正保证 worker 落在同一个 NVLink/RoCE 拓扑域里的，是 KAI topology annotation、ComputeDomain/DRA、以及 segmentSize 这类 K8s/infra 层约束。

特性分支里则不一样。至少这些 fetched refs 里已经有 topology-aware placement 逻辑：

1
2
3
4
5
6
7
8

origin/dsv3-from-v0.6.0
origin/topology-base-v2
origin/ultra-v3
origin/youngeunk/topology-aware-placement
origin/youngeunk/topology-aware-placement-dco-fix
origin/youngeunk/numa-aware-binding
origin/pull-request/2612
origin/pull-request/2613

这些分支的 nemo_rl/distributed/virtual_cluster.py 里，确实会读取 Ray node resource 中的 nvlink_domain_* / topo_rank，并把它们用于 bundle 排序；部分逻辑还支持通过 node_resource_constraints 把 placement group bundle pin 到某个 NVLink domain。

所以更准确的说法是：

main 分支已经有 infra probe，但 Python 消费侧还没合进来；特性分支里已经有完整一些的 topology-aware rank placement。

资源调度 Infra

K8s/KAI/DRA：容器资源 clique

GB300 相关 infra YAML 里，worker pod 有这几个关键配置：

1
2
3
4
5
6
7
8
9
10
11

metadata:
  annotations:
    kai.scheduler/topology: gb300-topology
    kai.scheduler/topology-required-placement: nvidia.com/gpu.clique
spec:
  schedulerName: kai-scheduler
  resourceClaims:
    - name: compute-domain-channel
      resourceClaimTemplateName: compute-domain-ultra-pipeclean
    - name: roce-channel
      resourceClaimTemplateName: roce-ultra-pipeclean

这部分不是 rank ordering，而是 pod placement 的硬约束。它让 KAI scheduler 以 nvidia.com/gpu.clique 为 topology placement key，把一组 worker gang-schedule 到合适的 GPU clique / ComputeDomain 里。没有这层，pod 可能被放到不同 rack，后面再怎么排 rank，也救不回跨 rack 通信。

segmentSize 也在这一层发挥作用。例如某个 worker group 有 64 个 replicas，segmentSize: 16 会把它展开成 4 个 16-worker segment：

1
2
3
4
5
6

for i in range(num_segments):
    segment = copy.deepcopy(wg)
    segment["groupName"] = f"{base_name}-segment-{i}"
    segment["replicas"] = segment_size
    segment["minReplicas"] = segment_size
    segment["maxReplicas"] = segment_size

这样 KAI 更容易把每个 segment 放进一个完整 clique。它解决的是“调度单位太大，无法表达 domain/rack 约束”的问题。

Ray worker startup：注册 topology resource

GB300 infra 的 worker command 会在 ray start 前跑一段 topology probe：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

CLUSTER_UUID=$(nvidia-smi -q 2>/dev/null \
  | grep 'ClusterUUID' \
  | head -1 \
  | awk -F: '{print $2}' \
  | tr -d ' ')

TOPO_RANK=$(hostname -i | tr -d '.')

RAY_RESOURCES='{'
if [[ -n "$CLUSTER_UUID" ]]; then
  RAY_RESOURCES+="\"nvlink_domain_${CLUSTER_UUID}\": 1"
  [[ -n "$TOPO_RANK" ]] && RAY_RESOURCES+=", "
fi
if [[ -n "$TOPO_RANK" ]]; then
  RAY_RESOURCES+="\"topo_rank\": $TOPO_RANK"
fi
RAY_RESOURCES+='}'

eval "$KUBERAY_GEN_RAY_START_CMD --resources='$RAY_RESOURCES'"

它把 NVLink fabric 的 ClusterUUID 变成 Ray custom resource：

1
2

nvlink_domain_<ClusterUUID>: 1
topo_rank: <some number>

注意这里的资源注册本身并不等于“rank placement 已经对齐”。Ray custom resource 只有在后续 placement group bundle 或 actor/task .options(resources=...) 里被请求时，才会形成调度约束；或者被 Python 侧读取后参与排序时，才会影响 rank order。

这也是 main 分支和特性分支的分叉点。

Rank placement：main 与特性分支

main 分支：只按 node/GPU 稳定排序

当前 main 的 RayVirtualCluster._get_sorted_bundle_indices() 逻辑大致是：

1. 用 placement_group_table(pg) 拿到每个 bundle 落在哪个 Ray node：

1

bundle_to_node_ids = pg_data["bundles_to_node_id"]

2. 对每个 bundle 启一个很小的 actor，强制 schedule 到那个 bundle 上，读它看到的 GPU id：

1
2
3
4
5
6
7
8

GetGPUIDActor.options(
    num_cpus=0.01,
    num_gpus=0.01,
    scheduling_strategy=PlacementGroupSchedulingStrategy(
        placement_group=pg,
        placement_group_bundle_index=i,
    ),
).remote()

3. 得到：

1

bundle i -> node_id, gpu_id

4. 按 (node_id, gpu_id) 排序：

1
2
3
4
5
6
7
8
9

bundle_infos = [
    (i, bundle_to_node_ids[i], gpu_ids[i])
    for i in range(num_bundles)
]

pg_reordered_bundle_indices = [
    bundle_info[0]
    for bundle_info in sorted(bundle_infos, key=lambda x: (x[1], x[2]))
]

为什么要做这个？因为 Ray placement group 里的 bundle index 不一定天然就是：

1
2
3
4
5
6
7

node0 GPU0
node0 GPU1
node0 GPU2
node0 GPU3
node1 GPU0
node1 GPU1
...

它可能是乱序的，比如：

1
2
3
4
5
6
7
8

bundle 0 -> nodeB GPU1
bundle 1 -> nodeA GPU2
bundle 2 -> nodeB GPU0
bundle 3 -> nodeA GPU0
bundle 4 -> nodeA GPU1
bundle 5 -> nodeB GPU2
bundle 6 -> nodeA GPU3
bundle 7 -> nodeB GPU3

如果直接按 bundle index 创建 rank，就会变成：

1
2
3
4
5

rank0: nodeB GPU1
rank1: nodeA GPU2
rank2: nodeB GPU0
rank3: nodeA GPU0
...

这对 TP/PP/DP/EP locality 都不友好。

排序后得到：

1
2
3
4
5
6
7
8

bundle 3 -> nodeA GPU0
bundle 4 -> nodeA GPU1
bundle 1 -> nodeA GPU2
bundle 6 -> nodeA GPU3
bundle 2 -> nodeB GPU0
bundle 0 -> nodeB GPU1
bundle 5 -> nodeB GPU2
bundle 7 -> nodeB GPU3

即：

1

_sorted_bundle_indices = [3, 4, 1, 6, 2, 0, 5, 7]

policy 侧再用这个排序后的 bundle list 创建 worker：

1
2
3
4
5

group_size = cluster.num_gpus_per_node
tied_groups = [
    (i // group_size, [bundle_idx])
    for i, bundle_idx in enumerate(cluster._sorted_bundle_indices)
]

如果 group_size = 4，就得到：

1
2

logical node 0: bundle 3,4,1,6  # actual nodeA GPU0..3
logical node 1: bundle 2,0,5,7  # actual nodeB GPU0..3

最后 RayWorkerGroup 按这个 bundle_indices_list 创建 worker，并设置：

1
2
3
4

RANK       = global_rank
LOCAL_RANK = bundle_idx
WORLD_SIZE = self.world_size
NODE_RANK  = pg_idx

所以 main 分支里的第 4 步，本质是：

在已经由 K8s/KAI 放好的资源上，把 Ray bundle 顺序整理成稳定的 node-major、gpu-major rank 顺序。

它没有读 nvlink_domain_*。

特性分支：真正消费 nvlink_domain_*

特性分支上的实现更接近最初想象中的“对齐 NVLink domain”。

以 origin/dsv3-from-v0.6.0 为例，virtual_cluster.py 里多了这些常量：

1
2
3
4

NVLINK_DOMAIN_PREFIX = "nvlink_domain_"
TOPO_RANK_KEY = "topo_rank"
NVLINK_DOMAIN_UNKNOWN = "unknown"
TOPO_RANK_UNKNOWN = -1

还多了一个 _get_gpu_id_info()：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

@ray.remote(num_gpus=1)
def _get_gpu_id_info() -> tuple[int, str, int]:
    gpu_id = ray.get_gpu_ids()[0]
    nvlink_domain = NVLINK_DOMAIN_UNKNOWN
    topo_rank = TOPO_RANK_UNKNOWN

    runtime_ctx = ray.get_runtime_context()
    node_id = runtime_ctx.get_node_id()
    for node in ray.nodes():
        if node.get("NodeID") == node_id:
            all_node_resources = node.get("Resources", {})
            break

    for key, val in all_node_resources.items():
        if key.startswith(NVLINK_DOMAIN_PREFIX):
            nvlink_domain = key
        if key == TOPO_RANK_KEY:
            topo_rank = int(val)

    return gpu_id, nvlink_domain, topo_rank

于是每个 bundle 的信息从 main 分支的：

1

(bundle_idx, node_id, gpu_id)

升级成：

1

(bundle_idx, node_id, gpu_id, nvlink_domain, topo_rank)

排序也从：

1

(node_id, gpu_id)

变成：

1

(domain_min_topo_rank, topo_rank, gpu_id)

对应代码里的描述是：

1
2
3
4

if topology_info_available:
    sort_key = (domain_min_topo_rank, topo_rank, gpu_id)
else:
    sort_key = (node_id, gpu_id)

这就真的把 nvlink_domain_* 纳入 rank ordering 了。

可选能力：domain pinning

更进一步，特性分支的 RayVirtualCluster.__init__() 多了：

1
2

segment_size: int | None = None
node_resource_constraints: list[dict[str, float]] | None = None

node_resource_constraints 的注释很直白：

1
2
3

node_resource_constraints = [
    {"nvlink_domain_<uuid>": 0.001},
] * 16

这表示给 16 个 logical nodes 都追加同一个 Ray custom resource requirement，从而把它们 pin 到同一个物理 NVLink domain。

创建 placement group bundle 时：

1
2
3
4
5

def _make_bundle(node_idx: int) -> dict:
    bundle = {"CPU": num_cpus_per_bundle, "GPU": num_gpus_per_bundle}
    if self.node_resource_constraints and self.node_resource_constraints[node_idx]:
        bundle.update(self.node_resource_constraints[node_idx])
    return bundle

这时 nvlink_domain_* 不只是“排序信息”，而是可以进入 Ray placement group 的 resource request，变成调度约束：

1

{"CPU": 1, "GPU": 1, "nvlink_domain_<uuid>": 0.001}

这和 main 分支有本质区别。main 只注册了资源，但 placement group bundle 没请求它；特性分支可以请求它。

vLLM：按 topology-sorted bundles 切 TP/PP group

特性分支的 vllm_generation.py 也相应改了。

当 vLLM 需要跨节点 model parallelism 时，它会用 unified placement group。之后 _get_tied_worker_bundle_indices() 会调用内部的 allocate_worker_groups()，传入：

1
2

sorted_bundle_indices=cluster._sorted_bundle_indices
nvlink_domain_per_bundle_index=cluster._nvlink_domain_per_bundle_index

如果有 topology-sorted bundle list，就直接按这个 list 切连续片段：

1
2

flat = list(sorted_bundle_indices)
slice_ = flat[i * model_parallel_size : (i + 1) * model_parallel_size]

每个 slice 是一个 model-parallel group，也就是一个 DP replica 里的 TP/PP workers。

如果一个 slice 里出现多个 NVLink domain，代码会 warning：

1
2
3

Model-parallel group ... spans ... NVLink domains;
cross-domain collectives may use slower links.
Prefer TP*PP that divides usable GPUs per domain.

这点很重要：排序能让连续 rank 尽量落在同一个 domain，但如果 TP * PP 和每个 domain 的可用 GPU 数不整除，还是可能切出跨 domain 的 group。拓扑感知不是魔法，parallelism 配置还得和物理 domain 尺寸匹配。

Megatron/DeepEP：HybridEP 是另一层

NeMo-RL 的 Megatron setup 里还有：

1
2

NUM_OF_HYBRID_EP_RANKS_PER_NVLINK_DOMAIN
USE_MNNVL

这部分是给 DeepEP/Megatron HybridEP 用的。

如果 config 里有：

1
2
3
4

megatron_cfg:
  moe_flex_dispatcher_backend: hybridep
  hybridep_num_ranks_per_nvlink_domain: 72
  hybridep_use_mnnvl: true

setup 会写：

1
2

os.environ["NUM_OF_HYBRID_EP_RANKS_PER_NVLINK_DOMAIN"] = "72"
os.environ["USE_MNNVL"] = "1"

如果不显式设置，默认大致是：

1
2

NUM_OF_HYBRID_EP_RANKS_PER_NVLINK_DOMAIN = min(expert_model_parallel_size, 64)
USE_MNNVL = int(expert_model_parallel_size > 4)

这和 Ray custom resource nvlink_domain_* 不是一回事：

• nvlink_domain_*：Ray / placement 层看到的节点拓扑资源。
• NUM_OF_HYBRID_EP_RANKS_PER_NVLINK_DOMAIN：Megatron/DeepEP 运行时理解 HybridEP domain 大小的参数。

前者决定 rank 怎么被放；后者告诉通信/dispatcher 这些 rank 应该按多大的 NVLink domain 处理。理想情况下，二者要一致，否则上层以为的 domain 和实际 placement 可能对不上。

可视化

完整可视化放在独立页面里，避免大段 HTML 干扰正文渲染：

打开 NeMo-RL NVLink domain placement 可视化

小结

这次阅读最大的收获，是不要把“拓扑感知”当成一个单点功能。它横跨好几层，每层能保证的东西不一样：

• K8s/KAI/DRA 负责 资源真的落在正确的物理 domain。
• Ray custom resource 负责 把物理 domain 暴露给 Ray 调度/运行时。
• Ray placement group 负责 把 bundle 绑定到资源约束上。
• NeMo-RL RayVirtualCluster 负责 把 bundle 顺序转成 rank 顺序。
• Megatron/DeepEP 负责 按这些 rank 做 TP/PP/EP/MoE 通信。

当前 main 只完成了其中一部分：infra 注册了 nvlink_domain_*，但 Python rank placement 没消费它。特性分支补上了后半截：读 nvlink_domain_* / topo_rank，按 topology 排 bundle，必要时通过 node_resource_constraints 把 placement group bundle pin 到 domain。

所以如果问：

NeMo-RL 现在是不是用 nvlink_domain_* 对齐 rank placement？

答案要分支限定：

• main：没有。main 主要靠 K8s/KAI 保证 placement，NeMo-RL 只做 (node_id, gpu_id) 稳定排序。
• topology-aware 特性分支：有。它已经把 nvlink_domain_* 纳入排序和可选调度约束。

这个差异也解释了为什么 main 的 infra 注释里已经写着“Registers nvlink_domain_<UUID> and topo_rank so virtual_cluster.py can do segment-aware rank assignment”，但实际 main 代码里搜不到消费路径：看起来 Python 侧实现还停在特性分支/PR 上，没有合入主线。

Nemotron 3 Ultra 的技术报告标题是 Nemotron 3 Ultra: Open, Efficient Mixture-of-Experts Hybrid Mamba-Transformer Model for Agentic Reasoning。

报告概览

报告围绕四大板块展开：

1. 预训练（Pretraining）
   架构细节（层数、MoE 配置、Mamba 头参数等）、NVFP4 低精度训练配方、数据混合策略（两阶段课程：Phase 1 偏多样性、Phase 2 偏质量）、长上下文扩展（用 32-way context parallelism 在 GB200 上扩展到 1M），以及两次训练发散事件的诊断（原因、回滚方案）。
2. 后训练（Post-training）
   SFT → RLVR → MOPD（多教师蒸馏）的三阶段流水线。训练了超过 10 个专项教师模型，涵盖软件工程、终端使用、搜索、Office 办公任务、安全、数学推理、竞赛编程等领域。还包括 MTP Boosting（推测解码头的专项优化）和推理预算控制。
3. 量化（Quantization）
   NVFP4 量化方案及其对精度的影响分析。
4. 推理（Inference）
   吞吐量对比、MTP 加速推理的效果数据。

RL Infra 总览

报告的 Infra 内容集中在 RL/MOPD 训练阶段的工程挑战，背景是 GB200 NVL72 集群 + Slurm + Ray + vLLM + NeMo-RL 的异构分布式训练栈。以下按问题分类整理：

这部分最有价值的地方，是它没有只给一个大而化之的“我们用了很多 GPU”描述，而是把 post-training 期间真实会拖慢训练、放大失败、拖垮共享服务的问题列出来，并给出前后对比。对 RL infra 来说，核心目标不只是单个 GPU 的 token/s，而是让 policy update、rollout generation、环境执行、judge / reward、checkpoint 和重启恢复共同组成一个稳定的 training loop。

报告里 RL/MOPD 采用的是 one-step off-policy asynchronous RL：rollout generation 和 policy update 会 overlap，整体 step time 取决于更慢的一侧。NVIDIA 观察到瓶颈通常在 rollout generation，而 rollout 又被少量长尾 generation 拖慢。所以后面的 MTP、vLLM、多节点启动、JIT cache、I/O cache，本质上都是为了减少 RL step 的尾延迟和重启成本。

从失败归因看，报告把 RL 软件失败分成三类：

这说明大规模 agentic RL 的主要风险已经不只是 trainer 本身，而是 生成引擎 + 环境 / 工具执行系统。二者合计约 92% 的失败，意味着 infra 要把 vLLM、Ray actor、sandbox、tool calling、health check、timeout 和 cleanup 都纳入训练系统，而不是把它们当作外部黑盒。

关键优化

MTP 加速 Rollout 生成

RL/MOPD 训练中，rollout 生成是瓶颈，而生成时间又被极少数"straggler"长请求主导。方案是用 MTP 投机解码（MTP head 先出 k 个候选 token，base model 单次前向验证），k=5 时带来 1.46× 加速，收益集中在长尾慢请求上。

这里的关键点是，MTP 在 infra 里的作用不是单纯提高在线推理吞吐，而是缩短训练循环里的 rollout wall-clock。长尾请求往往会在 batch 后段以较低并发继续 decode，投机解码在这类场景里收益更明显。也就是说，它优化的是 RL step 的 tail latency，而不是平均单请求延迟。

Slurm 启动优化

原有启动脚本对每个节点发多次独立 srun 来分配角色，slurmctld 串行处理 RPC，节点数多时复杂度 O(n)，导致启动耗时 30+ 分钟。改为单次多节点 srun，Slurm 控制器交互降为 O(1)，启动时间从 30+ 分钟降到 10 分钟。

这个问题很典型：小规模下“每个节点单独发命令”很直观，但到几百、上千节点后，控制面会先于数据面变成瓶颈。这里的优化不是更快的 GPU kernel，而是减少对 Slurm controller 的 RPC 压力，让 job launch 从节点数线性增长回到常数级控制面交互。

Ray GCS 扩展性

3000+ GPU 规模下，RL job 同时创建大量 Ray actor（policy、generation、env、judge 四类 worker），Ray 单线程 GCS 被 actor 注册请求淹没，启动耗时 25–49 分钟，且出现雷群式启动失败。解法：

• 把短生命周期 actor 改为 task，减少 40% actor 注册量
• 每节点合并初始化 actor
• 升级到 Ray 2.55（Anyscale 在此版本中修了 GCS 扩展性回归）

这里反映出 Ray 在 RL infra 里的角色更像一个异构 job runtime，而不是只负责几个 Python worker。policy、generation、environment、judge 的生命周期和资源需求不同，如果全部抽象成大量 actor，会把 GCS 注册路径打爆。短生命周期逻辑改为 task，本质是降低控制面对象数量；每节点合并初始化 actor，则是减少 thundering herd。

这也说明 RL infra 的调度单位要谨慎设计：不是所有并发实体都适合作为独立 actor，尤其是在数千 GPU、数千到数万进程同时启动的场景下，控制面元数据本身就是稀缺资源。

NVLink 拓扑感知调度

GB200 NVL72 的 NVLink 域覆盖单个 rack（18 节点 × 4 GPU = 72 GPU）。若 Megatron 的 Expert Parallelism (EP) 组跨 rack，MoE all-to-all 通信就会走 InfiniBand 而非 NVLink。原因是 Ray node ID 是随机 UUID，无拓扑信息，Megatron 按收到的 rank 排序分组，无法感知物理位置。

解法三步走：

1. 容器启动时，用 nvidia-smi -q 读取 NVLink fabric ClusterUUID，注册为 Ray 自定义资源，给调度器提供 rack 归属信息
2. NeMo RL 的 RayVirtualCluster 按 (domain_min_topo_rank, topo_rank, gpu_id) 组合键排序 bundle，确保 EP 组内所有 GPU 在同一 NVLink 域；用 SLURM_TOPOLOGY_ADDR 修正 Slurm block 乱序
3. Megatron 启用 external_gpu_device_mapping=True，信任 Ray 的 GPU 绑定

效果：端到端吞吐提升 20%。

这部分是整篇 infra 里最值得关注的优化。GB200 NVL72 的 rack 内 72 GPU 通过 NVLink 域连接，MoE 的 EP all-to-all 如果落在同一个 NVLink 域内，通信路径和跨 rack InfiniBand 完全不是一个量级。问题在于 Ray 和 Megatron 默认看到的是逻辑 rank，而不是物理 rack / NVLink domain。

因此这里其实做了一个跨层契约：

这个思路比“给训练框架加一个 topology 参数”更完整：它把硬件拓扑、Slurm 拓扑、Ray placement group、Megatron rank assignment 串成了一条链。链上任意一层丢失物理信息，EP 组就可能静默跨 rack，吞吐下降但不一定立刻报错。

NUMA 绑定

GB200 NVL72 每个计算托盘有一块 Grace CPU，含两个 socket 和多个 NUMA 节点：GPU 0/1 对应 NUMA node 0，GPU 2/3 对应 NUMA node 1。NVLink-C2C 是 socket-local 的，若 Ray worker 进程跑到远端 socket，CPU↔GPU 的内存带宽会因跨 socket coherence link 而下降。

解法：显式把 policy worker 和 vLLM worker 绑定到与其分配 GPU 物理相邻的 CPU socket。优化器 state offload、tokenization、数据预处理、pinned-memory 申请全部命中本地 DRAM 和 C2C 路径。效果：端到端吞吐提升 10%。

这说明在 Grace-Blackwell 这类 CPU/GPU 紧耦合平台上，CPU placement 已经会影响端到端训练吞吐。对 RL 来说，CPU 不是闲置资源：tokenization、环境执行、judge、数据预处理、checkpoint D2H staging、vLLM worker 管理都会用 CPU。如果进程被 Linux 调度到远端 socket，就会把原本应该走本地 C2C 的路径变成跨 socket 访问。

Checkpoint 异步保存

同步 checkpoint 阻塞训练约 60 秒/次。优化链路：

• 启用 NVRx（NVIDIA Resiliency Extension）异步 checkpoint：参数先 copy 到 CPU，后台持久化，阻塞降到 6–8 秒
• NCCL D2H 传输与 CPU 拷贝 overlap，缩短同步窗口
• 持久化 checkpoint worker 进程，避免每次 fork/spawn 开销
• checkpoint finalization（跨 worker 同步 + 最终 commit）移入后台线程，完全不阻塞训练
• 缓存 distributed save plan，不再每次重新计算

最终效果：exposed save time 降到 <1 秒。

这个 checkpoint 优化有两层含义。第一层是减少训练停顿：把参数拷到 CPU 后后台落盘，让 GPU 训练尽快继续。第二层是把 checkpoint 的分布式协议开销从 hot path 移出去：finalization、save plan 计算、worker 进程创建都不能每次阻塞训练主循环。

Megatron Core Distributed Optimizer 也很关键，因为 optimizer state 已经按 data-parallel rank 分片，每个 rank 只保存本地 shard。对超大模型 post-training 来说，checkpoint 不只是容灾机制，也是影响 goodput 的常规周期性开销。

JIT 编译缓存

1000+ GPU 冷启动时，初始化耗时 ~49 分钟，其中 JIT 编译占 38.8 分钟，明细如下：

解法三步：

1. 共享存储持久化缓存：job 结束时把所有 JIT artifacts 压缩打包写回共享目录
2. 节点本地 seeding：Ray 初始化前，每个节点把压缩包解压到本地 /tmp，后续 JIT 写也打到 /tmp，避免并发元数据风暴
3. FlashInfer cubin 预编译进容器镜像：build 时用 flashinfer download-cubin bake into image，消除运行时编译

效果：初始化时间从 38.8 分钟降到 0.4 分钟，降低 99%。

JIT cache 这一节可以理解为“把运行时不确定性前移或局部化”。FlashInfer cubin、Inductor graph、Triton autotuning、vLLM CUDA graph capture 都会生成本地 artifact；如果每个 worker 冷启动都重新编译，集群规模越大，重复工作越多，启动时间也越不可预测。

这里的三个手段分别解决不同问题：

值得注意的是，缓存不是简单放共享盘就结束了。千卡规模下，共享盘上的大量小文件读写会变成 metadata storm，所以训练时写本地、启动时顺序解压 tarball、结束时单点归档，才是完整方案。

多节点 vLLM 稳定性

多节点 vLLM 在 GB200 上遇到三类问题及对应修复：

此外还加了 vLLM health check、RPC timeout、graceful shutdown 和孤儿进程清理，提升整体稳定性。

这部分的工程信号很强：vLLM 在 RL 训练里不是一个独立 serving 服务，而是 Ray 分布式执行图里的 generation backend。主 Ray actor 还会通过 multiprocessing.spawn 拉起 EngineCore 子进程，子进程再连 GCS。这样一来，版本、环境变量、NCCL 初始化、CUDA graph、FlashInfer kernel 都必须在父进程、子进程、多节点之间保持一致。

所以它的稳定性问题往往不是“vLLM 算错了”，而是：

• 多组件安装的 GPU kernel 库版本不一致，导致 ABI 静默不兼容；
• 父进程的 LD_LIBRARY_PATH、CUDA/NCCL/FlashInfer 相关环境变量没有传到 spawn 子进程；
• 某些 JIT kernel 和 NCCL 多节点 NVLink memory registration 组合会 hang；
• 分布式 worker 没有 timeout / health check，单点卡死会拖住整个 RL job。

这些问题都很像生产 serving，但影响的是训练 goodput。RL 训练里的 generation backend 如果没有 fail-fast 和 cleanup，失败会从一个子进程扩散到整轮 rollout，最后变成整 job 重启。

容器镜像 I/O 风暴

44GB squashfs 容器镜像在千卡规模下被所有节点并发读取，共享存储被打爆，部分节点报 I/O error 或卡 12+ 分钟（正常应 2–3 分钟），一个慢节点就能拖延整个 job 的 Ray 初始化。job 结束时所有节点并发回写 JIT 缓存，制造第二次 I/O 风暴。

解法：

• Enroot 本地 squashfs 缓存：节点首次提取后，后续 job 直接复用本地缓存，不再读共享存储
• 写入非对称化：训练时 JIT 全写本地，job 结束时仅由一个 sidecar 进程归档压缩包回共享存储（因所有节点编译的 kernel 完全相同，只需保存一份）

这里可以和 JIT cache 放在一起看：大规模训练的启动阶段会同时制造两类 I/O 压力，一类是所有节点读同一个 44GB squashfs 镜像，另一类是所有节点写回大量 JIT 小文件。前者是吞吐型读风暴，后者是 metadata / 小文件写风暴。

NVIDIA 的处理方式都遵循同一个原则：共享存储只做低频、顺序、少并发的持久化；高频读写放节点本地。容器镜像靠 Enroot 本地 cache 变成 warm start，JIT artifact 靠 sidecar 只保存一份，因为所有节点编译出的 kernel 相同，没有必要让每个节点都回写。

后续方向：更细粒度的容错

报告的 Future Work 指向两个主要方向：

1. fail-fast fault isolation：generation worker 或 sandbox 实例出问题时，尽快隔离失败，避免 retry 引发级联故障。
2. component-level recovery：单个 generation worker、sandbox、tool-calling 组件可以独立重启，不必重启整个 RL job。
3. 细粒度状态 checkpoint：保存 in-flight rollout、KV cache、conversation state，让恢复从最近一致快照继续，而不是从头 replay。

这说明下一阶段 RL infra 的重点会从“让大 job 跑起来”转向“让局部失败不要扩大成全局失败”。尤其是 agentic RL 里，环境、工具、浏览器、终端、judge、模型生成都可能失败，训练系统需要像分布式在线系统一样做 fault domain 切分。

开销与收益

这份 infra 章节的工程密度相当高，尤其是 NVLink 拓扑感知调度和 JIT 缓存部分，对做大规模 RL 训练基础设施的团队很有参考价值。

RL Infra 启发

我的理解是，Nemotron 3 Ultra 的 infra 经验可以抽象成几条原则：

1. RL goodput 优先于单点吞吐：rollout 长尾、checkpoint、启动、恢复、工具执行失败都会进入 wall-clock。
2. 控制面要按规模重新设计：Slurm RPC、Ray GCS actor 注册、Ray actor spawn、vLLM 子进程连接都会在千卡规模暴露瓶颈。
3. 拓扑信息必须跨层传递：硬件 NVLink domain、Slurm topology、Ray placement、Megatron rank assignment 必须对齐。
4. 共享存储不能承受所有节点同时读写：镜像、JIT cache、checkpoint 都要区分本地热路径和共享持久化路径。
5. agentic RL 的 fault domain 更复杂：generation engine、sandbox、tool calling、judge、environment 都是训练系统的一部分，需要 health check、timeout、cleanup 和局部恢复。

这也是它和很多只讲模型结构的技术报告不同的地方：§3.6 把 RL 训练看成一个包含推理引擎、调度系统、容器系统、存储系统和任务环境的生产级分布式系统。对做训练平台的人来说，这部分比 benchmark 表格更值得反复看。

一个构想：在 Kubernetes 之上构建云原生分布式 RL 系统

如果把 Nemotron 3 Ultra 这套经验迁移到云原生体系里，一个自然的方向是 在 Kubernetes 之上构建云原生分布式 RL 系统。这里的 Kubernetes 不是 RL 系统本身，而是资源抽象、调度、隔离、生命周期管理和控制面的底座；真正的 RL 系统由 Ray / NeMo-RL / vLLM、learner、rollout generator、sandbox、tool calling、judge、checkpoint、cache 和 topology-aware runtime 共同组成。

所以这套系统不是把 Ray job 简单塞进 Pod，也不是只做一个训练 launcher，而是把 RL 训练里的 policy、rollout、inference engine、sandbox、tool calling、judge、checkpoint、cache 和 topology placement 都变成 Kubernetes 可理解、可调度、可恢复的资源。

这里有一个很关键的现实基础：NeMo-RL 本身已经开源，而且采用 Apache-2.0 License。官方 README 明确说它是 open-source post-training library，支持从小规模实验到 multi-GPU / multi-node deployment，并用 Ray 做资源管理和调度；GitHub repo 也标注了 Apache-2.0 License。也就是说，它不是只公开了 paper 或 recipe，而是把 RL post-training runtime 和相当一部分 infra 代码也放了出来。

公开仓库里能看到几类和 §3.6 对应的基础设施代码，而且不少已经不是“示例级说明”，而是直接在 NeMo-RL runtime 路径里：

所以这件事比“未来可以做一个 Kubernetes RL 系统”更现实：NeMo-RL 已经开源了 RL runtime、Ray 编排、Slurm/Kubernetes 启动、checkpoint 管理、NVLink topology plumbing、NUMA affinity 输入和 vLLM generation 稳定性 patch。剩下真正值得平台化的是，把这些能力从 repo 里的脚本、runtime 模块和 manifests，进一步沉淀成云原生控制器、调度插件、节点 daemon、cache controller 和标准化 CRD。

这套系统可以分成几层：

核心设计目标应该是 RL goodput first。Kubernetes 不能只知道“Pod Running”，还要知道这一组 Pod 是否真的在产生有效 rollout、是否有 generation tail、是否有 sandbox timeout、是否有某个 vLLM worker 卡死、是否因为 checkpoint 或 JIT cache 拖慢训练。换句话说，RLJob 的状态机要比普通 Job 更接近一个长生命周期分布式系统。

调度层可以重点做三件事：

1. 拓扑感知放置：把 GPU UUID、NVLink ClusterUUID、NUMA node、IB leaf / spine、rack 信息注册成 Node label / extended resource / scheduler plugin 输入，让 TP/EP 组优先落在同一 NVLink domain，跨域通信显式可见。
2. 异构 worker 编排：learner、rollout generator、sandbox、judge 的资源形态不同，不能只按 GPU 数量调度；CPU、memory、local NVMe、network、ephemeral storage 都会影响 RL step time。
3. 局部失败恢复：generation worker、sandbox、tool-calling worker 可以独立重启和替换，RLJob controller 负责隔离失败、保留 in-flight 状态、避免整 job 重启。

存储和启动路径也应该平台化。镜像、JIT cache、checkpoint 都走“本地热路径 + 共享持久化路径”：节点本地 NVMe 承担高频读写，共享存储只做少并发归档；checkpoint 异步化，save plan 缓存化；JIT artifact 可以通过 DaemonSet 预热到节点，或者在镜像构建阶段 bake into runtime layer。

这样的系统最终会像一个运行在 Kubernetes 之上的 RL 训练操作系统：Kubernetes 提供资源、隔离和控制面，Ray / NeMo-RL / vLLM 提供执行 runtime，平台层负责 topology、cache、checkpoint、health、failure domain 和 goodput。Nemotron 3 Ultra §3.6 的意义就在这里：它把未来 RL infra 要解决的问题摊开了，而云原生版本要做的，是把这些经验沉淀成可复用的控制器、调度插件和运行时协议。

本文以当前 NCCL master 分支源码为例，版本标识来自 makefiles/version.mk，当前为 NCCL 2.30.7。

这里讨论的“业务侧 NCCL clique”不是 NCCL 暴露的独立 API，而是指业务或训练框架对一组 rank / GPU 的通信组织方式。例如数据并行组、张量并行组、流水并行组、专家并行组，在运行时会对应到一个或多个 NCCL communicator。

相关阅读：

• NCCL 从建链到通信：解释 communicator 初始化、建链和一次通信执行主线。
• reading nccl：以 NCCL 2.14.3 为例展开 net_ib、RDMA、QP、MR、FIFO 等实现细节。
• NCCL clique 可视化页面：用交互图解释训练框架、launcher / rendezvous、调度 placement 和 NCCL communicator 的关系。

概念边界

三类 clique

NCCL 相关讨论里，clique 可以对应三类对象。

本文把业务文档中的“NCCL clique”对应到第一类：一组业务 rank 被组织成同一个 communicator。MNNVL clique 后文作为硬件 fabric domain 对照展开。

下图给出从业务输入到 NCCL 实际通信顺序的层次关系：

业务通信组到 communicator 的映射

NCCL API 注释里，ncclCommInitRank 要求 rank 在一个 communicator clique 内唯一。这里的 communicator clique 表示一组共同参与同一个 communicator 的 ranks。

典型初始化输入是：

1
2
3
4

ncclUniqueId
nranks
rank
CUDA device

其中：

• ncclUniqueId 决定哪些进程 / 线程加入同一个 communicator。
• nranks 决定 communicator 的规模。
• rank 是这个 communicator 内的逻辑编号，必须唯一。
• 当前 CUDA device 决定该 rank 绑定到哪张 GPU。

单进程多卡时，ncclCommInitAll 是一个便利封装。源码注释也说明 devlist 的顺序定义了 communicator 内 processor 的 user-order。它内部会生成一个 ncclUniqueId，再用 ncclGroupStartInternal / ncclGroupEndInternal 批量初始化多张 GPU 上的 communicator。

多进程训练框架通常先完成 rendezvous，然后为不同通信组构造不同的 communicator。例如：

1
2
3
4
5
6
7
8
9
10
11
12

global group:
  ranks = [0, 1, 2, 3, 4, 5, 6, 7]

tensor parallel group:
  ranks = [0, 1, 2, 3]
  ranks = [4, 5, 6, 7]

data parallel group:
  ranks = [0, 4]
  ranks = [1, 5]
  ranks = [2, 6]
  ranks = [3, 7]

这些业务组如果使用 NCCL backend，会对应到不同的 NCCL communicator。每个 communicator 都有自己的 rank 空间。全局 rank 4 在某个 data parallel communicator 里可能是 rank 1；在 global communicator 里仍然是 rank 4。

业务侧输入项

业务侧对 NCCL 的影响，主要发生在初始化之前和 collective enqueue 之前。

并行策略如何展开成 rank group

训练框架通常先描述并行维度，而不是直接手写所有 NCCL communicator。一个常见的描述方式类似：

1
2
3
4
5

world_size = 16
tensor_model_parallel_size = 4
pipeline_model_parallel_size = 2
data_parallel_size = world_size / tensor_model_parallel_size / pipeline_model_parallel_size
                   = 2

这类配置表达的是 group pattern：

1
2
3

TP: 每 4 个 rank 组成一个 tensor parallel group
PP: 每条 pipeline 包含 2 个 stage
DP: 相同 TP/PP 位置上的 rank 组成 data parallel group

这里需要拆开四个职责：

• 训练框架解释并行策略，定义 (dp_rank, pp_rank, tp_rank) 这样的逻辑坐标，并通过 rank layout 把逻辑坐标映射到 global rank。
• launcher / rendezvous 生成或分配 global / local rank。例如 torchrun、MPI、Kubernetes training operator 或框架自带 launcher，会为每个训练进程提供 RANK、WORLD_SIZE、LOCAL_RANK 等信息。
• 训练调度系统负责资源 placement，选择节点、GPU、NIC 或 fabric domain，并影响 rank 到物理 GPU 的落点。
• NCCL 接收 communicator 成员、rank 编号、device / busId / fabric 信息等输入，并在这些输入上生成实际通信图。

因此，如果“并行策略”只指 tp_size、pp_size、dp_size，它还不能唯一决定 global rank id。如果框架同时确定了 rank layout，并且 launcher / rendezvous 已经给进程分配 global rank，具体 rank id group 就是确定的。

假设框架采用如下 rank layout：

1

rank = ((dp_rank * pp_size + pp_rank) * tp_size) + tp_rank

则 16 个 rank 可以展开成：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

TP groups:
  dp0, pp0: [0, 1, 2, 3]
  dp0, pp1: [4, 5, 6, 7]
  dp1, pp0: [8, 9, 10, 11]
  dp1, pp1: [12, 13, 14, 15]

PP groups:
  dp0, tp0: [0, 4]
  dp0, tp1: [1, 5]
  dp0, tp2: [2, 6]
  dp0, tp3: [3, 7]
  dp1, tp0: [8, 12]
  dp1, tp1: [9, 13]
  dp1, tp2: [10, 14]
  dp1, tp3: [11, 15]

DP groups:
  pp0, tp0: [0, 8]
  pp0, tp1: [1, 9]
  pp0, tp2: [2, 10]
  pp0, tp3: [3, 11]
  pp1, tp0: [4, 12]
  pp1, tp1: [5, 13]
  pp1, tp2: [6, 14]
  pp1, tp3: [7, 15]

如果框架换一种 rank layout，虽然 tp_size=4、pp_size=2、dp_size=2 没变，rank id group 也可能改变。例如：

1

rank = ((pp_rank * dp_size + dp_rank) * tp_size) + tp_rank

这时 rank 4 从前一个例子里的 (dp0, pp1, tp0) 变成 (dp1, pp0, tp0)。因此：

1
2
3
4
5
6
7

前一个 layout:
  DP group, pp0, tp0: [0, 8]
  PP group, dp0, tp0: [0, 4]

新的 layout:
  DP group, pp0, tp0: [0, 4]
  PP group, dp0, tp0: [0, 8]

这个对比说明：parallel size 决定逻辑 group pattern，rank layout 决定具体 rank id group，placement 决定这些 rank 的物理含义。即使框架已经生成 [0, 1, 2, 3] 这样的固定 group，把它放在一台 8 卡 NVSwitch 机器上，和分散在 4 台机器上，对 NCCL 来说仍然是不同的拓扑输入。

rank 集合

业务可以决定哪些进程 / GPU 进入同一个 communicator。

这会直接影响：

• communicator 的 nranks；
• collective 的参与者；
• NCCL 拓扑搜索看到的 rank 集合；
• 后续 ring、tree、NVLS、CollNet 等图的候选空间。

如果业务把跨机 rank 放进同一个 communicator，NCCL 需要考虑 NET transport。如果业务把同机 GPU 放进一个 communicator，NCCL 可以在 P2P、SHM、NVLS 等路径中选择。

rank 顺序

业务可以决定 communicator 内 rank 的编号顺序。

这个顺序是 NCCL 图搜索的输入之一，但不等价于“业务指定 ring 顺序”。NCCL 还会结合 GPU busId、NVLink、PCIe、NIC、NUMA、MNNVL 等拓扑信息计算 ring / tree。

rank 到 GPU 的 placement

业务调度和框架启动方式决定 rank 绑定到哪张 GPU。

该映射会影响 NCCL 的 transport 选择和图搜索，因为 NCCL 会使用物理位置信息：

• 同一进程还是不同进程；
• 同一节点还是跨节点；
• GPU 之间是 NVLink、PCIe 还是需要经过 CPU / NIC；
• GPU 到 NIC 的距离；
• 是否属于同一个 MNNVL fabric domain。

例如，同样是 8 个 rank，如果 rank 0-7 都在同一台 8 卡机器上，和 rank 0-7 分散在 8 台机器上，NCCL 得到的图和 transport 选择会不同。

communicator 切分

业务可以通过框架里的 process group / communication group，或者直接使用 ncclCommSplit，把已有通信世界切成多个 communicator。

ncclCommSplit 的语义是：

• 相同 color 的 ranks 会进入同一个新 communicator；
• key 用来决定新 communicator 内 rank 顺序；
• NCCL_SPLIT_NOCOLOR 表示该 rank 不进入任何新组。

这种语义可以对应业务里的 data parallel、tensor parallel、expert parallel 分组。它影响的是 communicator 边界，而不是单次 collective 内部的 packet 传输顺序。

collective 调用和 group launch

业务可以影响 collective 的 enqueue 顺序：

• 哪个 stream 上调用；
• 同一个 stream 上 collective 的先后顺序；
• 是否用 ncclGroupStart / ncclGroupEnd 把多个 communicator 的操作聚合起来；
• 每次 collective 的数据规模、datatype、op 类型。

NCCL 内部 ncclGroupCommJoin 会把同一进程内 sibling communicators 放在相邻位置。这里的 sibling 由 intraComm0 判断。doLaunches 再按这些本地 clique 分批 prepare、launch、finish。

因此，业务调用顺序会影响 launch 组织方式；NCCL 也会在本地做聚合和排序，以满足多 GPU 同进程场景下的 launch 约束。

环境变量和配置

业务或平台也可以通过环境变量影响 NCCL 选择。

常见方向包括：

• NCCL_ALGO：约束 algorithm，例如 Ring、Tree、NVLS、CollNet 等。
• NCCL_PROTO：约束 protocol，例如 Simple、LL、LL128。
• NCCL_MIN_NCHANNELS / NCCL_MAX_NCHANNELS：影响 channel 数量边界。
• NCCL_MNNVL_ENABLE：控制是否启用 MNNVL 检测。
• NCCL_MNNVL_CLIQUE_ID：在支持 MNNVL 的系统上，用于设置 MNNVL cliqueId。
• NCCL_MNNVL_CROSS_CLIQUE：允许同一 NVLink domain 内跨 clique 的 P2P。

这些配置可用于验证、隔离或特定平台调优。它们约束 NCCL 的选择空间，不等同于显式指定完整的通信执行计划。

MNNVL clique

MNNVL clique 表示 Multi-Node NVLink fabric 中的硬件通信域。它和业务 clique 都描述 rank 集合，但来源是 NVML / Fabric Manager 提供的 fabric 信息。

fabric 信息来源

初始化时，NCCL 会在 fillInfo 里通过 NVML 读取 GPU fabric 信息。相关信息包括：

1
2
3

fabricInfo.clusterUuid
fabricInfo.cliqueId
fabricInfo.state

当前源码也提供了 NCCL_MNNVL_UUID 和 NCCL_MNNVL_CLIQUE_ID 这样的环境变量覆盖入口。文档里说明 NCCL_MNNVL_CLIQUE_ID 从 2.25 起可用，通常由 Fabric Manager 分配，也可以用于 MNNVL job 的 soft partition。

NCCL 判断逻辑

随后 ncclMnnvlCheck 会做几件事：

• 要求 cuMem 和 FABRIC handle 支持可用；
• 要求所有 ranks 的 fabric state 都完成；
• 比较 rank 之间的 clusterUuid；
• 在相同 clusterUuid 内，再按 cliqueId 判断是否属于同一个 clique；
• 把同 clique 的 ranks 记录到 comm->clique.ranks；
• 设置 comm->cliqueRank、comm->clique.size、comm->nvlDomainSize；
• 条件满足后设置 comm->MNNVL = 1。

拓扑检查里，ncclTopoCheckMNNVL 也是按这个规则判断：

1
2

same clusterUuid &&
  (same cliqueId || p2pCrossClique enabled)

如果启用了 cross-clique，并且两个 rank 在同一个 UUID 但不同 cliqueId 中，NCCL 会把它作为 cross-clique P2P 场景处理。

与业务 clique 的关系

二者关系可以描述为：

1
2
3
4
5
6
7
8
9

业务 clique:
  由业务决定哪些 rank 组成一个 communicator。

MNNVL clique:
  由 Fabric / NVML 信息描述哪些 rank 处在同一个 Multi-Node NVLink clique。

二者交集:
  如果业务把一组 rank 放进 communicator，而这些 rank 在硬件上也属于同一个 MNNVL clique，
  NCCL 可以把 MNNVL 纳入 topology / transport 选择。

通信顺序

“NCCL 的通信顺序”包含多个层次。业务输入的影响范围取决于具体层次。

communicator 内 rank 顺序

业务通过 rank list、key、devlist、进程启动顺序等方式，决定 communicator 内 rank 编号。NCCL 后续所有拓扑和 collective 逻辑都以 communicator 内 rank 编号为基础。

后续 topology graph 会在这个输入上继续计算。

topology graph 顺序

NCCL 初始化时会根据 rank 到 GPU / NIC 的物理关系计算图。

在 ring 场景下，最终每个 channel 上，本 rank 看到的是：

1
2

channel[c].ring.prev
channel[c].ring.next

在 tree 场景下，则会有 parent / children 等关系。这个阶段综合了业务 placement、硬件拓扑、NCCL 配置和搜索策略。

业务可以通过 placement 和 rank mapping 影响 topology graph，但 rank list 的顺序不等同于最终 ring 顺序。

group launch 顺序

如果业务用 ncclGroupStart / ncclGroupEnd 聚合多个 communicator 或多个 GPU 上的 collective，NCCL 会做本地 launch 排序。

当前源码中，ncclGroupCommJoin 会优先把具有相同 intraComm0 的 communicator 放在一起。注释里明确说这样做是为了保留用户程序顺序，同时保证 sibling communicators 连续出现。

doLaunches 的外层循环也按相同 intraComm0 形成的 clique 迭代。每个 clique 内先 prepare，再按 rounds launch，最后 finish。

这里的 clique 是 launch clique，不是 MNNVL clique。

kernel 内 chunk / step 顺序

数据移动发生在 NCCL kernel 和 proxy 线程里。

这一层由多个因素共同决定：

• collective 类型：AllReduce、AllGather、ReduceScatter、Broadcast 等；
• algorithm：Ring、Tree、NVLS、CollNet、PAT 等；
• protocol：Simple、LL、LL128；
• channel 数量；
• chunk / slice / step 划分；
• transport 是 P2P、SHM 还是 NET；
• NET 场景下 proxy 如何推进 send / recv。

业务可以通过 message size、算法 / 协议环境变量、placement 影响选择结果，但具体 packet 或 chunk 的顺序由 NCCL kernel 和 transport progress 逻辑决定。

示例

假设有 16 个 GPU，分布在两台 8 卡机器上：

1
2

node0: rank 0 1 2 3 4 5 6 7
node1: rank 8 9 10 11 12 13 14 15

业务可以组织两类通信组：

1
2
3
4
5
6
7

tensor parallel:
  [0,1,2,3,4,5,6,7]
  [8,9,10,11,12,13,14,15]

data parallel:
  [0,8], [1,9], [2,10], [3,11],
  [4,12], [5,13], [6,14], [7,15]

从业务层看，这是两类 clique configuration。

从 NCCL 层看，它们会对应到不同 communicator：

• tensor parallel communicator 主要覆盖同机 GPU 通信，NCCL 可能选择 P2P / SHM / NVLS 等路径；
• data parallel communicator 跨节点，NCCL 需要考虑 NET transport；
• 如果平台支持 MNNVL，并且这些 GPU 的 fabric UUID / cliqueId 满足条件，MNNVL 会进入拓扑判断；
• 如果业务改变 rank 到 GPU 的 placement，两个 communicator 的拓扑图和 transport 选择都可能改变。

这个例子中，业务没有直接指定“ring 是 0->1->2->3”。业务指定的是：哪些 rank 一起通信，以及这些 rank 放在哪些 GPU 上。

观察和验证

如果要验证业务 clique 到 NCCL 实际图的映射，可以从日志和 dump 文件入手。

常用方式：

1
2
3
4

NCCL_DEBUG=INFO
NCCL_DEBUG_SUBSYS=INIT,GRAPH,TUNING,NET,COLL
NCCL_TOPO_DUMP_FILE=/tmp/nccl-topo.xml
NCCL_GRAPH_DUMP_FILE=/tmp/nccl-graph.xml

如果只想让特定 rank dump，可以结合：

1
2

NCCL_TOPO_DUMP_FILE_RANK=0
NCCL_GRAPH_DUMP_FILE_RANK=0

可以重点观察几类信息：

• communicator 的 rank、nRanks、localRank、MNNVL；
• ring / tree 图里每个 channel 的邻居；
• algorithm / protocol 的选择；
• NET / P2P / SHM / NVLS 相关日志；
• MNNVL 的 cliqueId、cliqueSize、cliqueRank、nvlDomainSize。

总结和启示

下图用四个职责边界概括前文：训练框架、launcher / rendezvous、训练调度系统和 NCCL。重点不是严格时间线，而是输入归属。

归纳起来，业务 clique 是 communicator membership，MNNVL clique 是硬件 fabric domain。前者来自框架通信组，后者来自 Fabric / NVML 信息；业务可以影响 MNNVL 的使用条件，但不能把它当作普通业务分组 API。

对训练业务侧，关键是把 communication group、rank layout、rank 到 GPU 的映射、stream 使用和 collective 调用顺序固定下来。业务可以定义哪些 rank 一起通信，也可以约束 algorithm、protocol 或 MNNVL 行为，但不能直接指定每个 channel、chunk、packet 的执行顺序。

对训练调度系统，关键是把 topology-aware placement 做成可解释的契约。调度器可以消费并行策略、通信亲和性，或者已经展开的 rank group 约束，然后输出满足约束的 node / GPU / NIC placement。只分配 GPU 数量，而不记录节点边界、NIC 亲和性、NVLink / MNNVL domain，就很难稳定复现业务期望的 NCCL 行为。

对于支持 MNNVL 的平台，Fabric Manager 分配的 UUID、cliqueId、fabric state 也应该进入节点画像、准入判断和 placement 结果解释中，而不是只在容器启动后交给 NCCL 自行发现。

因此，训练框架、launcher / rendezvous 和调度系统之间需要明确通信拓扑契约。契约越清晰，越容易定位性能差异来自并行策略、rank layout、rank assignment、调度 placement、硬件 fabric 状态，还是 NCCL algorithm / transport 选择。

以当前 NCCL master 分支源码为例，版本标识来自 makefiles/version.mk，当前为 NCCL 2.30.7。本文梳理 communicator 初始化时如何从拓扑搜索结果走到 transport 连接建立，以及建链完成后一次通信如何复用这些连接。

这里说的“建链”可以拆成两层：

1. 图层：决定每个 channel 上，本 rank 与哪个 rank 通信。
2. transport 层：决定这条边具体用 P2P、SHM 还是 NET 连接，并把 device kernel 需要的 ncclConnInfo 填好。

建链完成后，通信阶段主要复用 ncclConnInfo，由 host 生成 work，由 device kernel 和必要的 proxy 线程推进数据传输。

相关阅读：reading nccl 主要记录 NCCL 2.14.3 中 net_ib、RDMA、QP、MR、FIFO 等网络传输细节。本文侧重从 communicator 建链到一次通信执行的主线，两篇可以结合阅读。

关于业务侧如何组织 communicator、rank / GPU placement 如何影响 NCCL clique、MNNVL clique 和通信顺序，可以参考：业务侧配置与 NCCL clique。

总览

阶段主线

主线可以压缩为：

1
2
3
4
5
6
7
8

topology / graph
  -> channel 逻辑邻居
  -> connectSend / connectRecv bitmask
  -> transport.setup 生成 ncclConnect
  -> bootstrap 交换 handle
  -> transport.connect 填 ncclConnInfo
  -> cudaMemcpyAsync 到 devPeers
  -> device kernel 可直接通信

下图按阶段和模块边界展示建链过程：

主要源码入口

对应的主要文件：

• src/init.cc
  • initTransportsRank 是 communicator 建链主线。
  • ncclTopoCompute 搜索 ring/tree/collnet/nvls graph。
  • ncclTopoPostset 后，comm->channels[c].ring/tree 已经知道逻辑邻居。
  • 非 runtime connect 时，在初始化阶段直接调用 ncclTransportRingConnect 和 ncclTransportTreeConnect。
• src/graph/connect.cc
  • ncclTopoPreset：根据本节点 intra 排列先生成本 rank 的 topoRanks。
  • ncclTopoPostset：聚合所有 rank 的 topoRanks，生成全局 ring/tree 关系。
• src/transport/generic.cc
  • ncclTransportRingConnect
  • ncclTransportTreeConnect
  • ncclTransportPatConnect
• src/transport.cc
  • ncclTransportP2pConnect：标记要连谁。
  • ncclTransportP2pSetup：执行 setup、交换 handle、connect。
  • selectTransport：按 transport 优先级选择连接方式。
• src/transport/p2p.cc
• src/transport/shm.cc
• src/transport/net.cc

与 reading nccl 的内容对照

reading nccl 可以看作本文中 NET/IB 传输部分的展开。本文描述 NCCL 的通用抽象和调用主线，reading nccl 更接近 ncclNetIb 的实现细节。

阅读时可以先用本文建立两条主线：

1
2

建链：topology -> transport -> ncclConnInfo
通信：work -> device primitive -> proxy -> ncclNet

再回到 reading nccl 看 IB verbs 细节：

1
2
3
4
5

ncclNet -> ncclNetIb
  -> connect / accept
  -> isend / irecv
  -> test
  -> ibv_post_send / ibv_post_recv / ibv_poll_cq

阅读路径

如果直接从 initTransportsRank 沿调用链展开，会同时涉及 topology、bootstrap、proxy、CUDA IPC、SHM、NET、GDR、PXN 等多个子系统。

阅读时可以先按问题边界分层，再进入完整调用链。

第一层：先看谁连谁

这一层只看 graph 层，不关心 P2P、SHM、NET。

要回答的问题是：

1

每个 channel 上，本 rank 的上游和下游分别是谁？

关注函数：

1
2
3
4
5
6
7
8
9
10
11
12

src/init.cc
  initTransportsRank
    ncclTopoCompute(ringGraph)
    ncclTopoCompute(treeGraph)
    ncclTopoPreset
    ncclTopoPostset

src/graph/connect.cc
  ncclTopoPreset
  ncclTopoPostset
  connectRings
  connectTrees

这一层的结论：

1
2
3
4
5
6

topo / graph 阶段只是在填逻辑邻居：

channel->ring.prev
channel->ring.next
channel->tree.up
channel->tree.down[]

也就是说，这一层还没有真的建立连接，只是把通信图画出来。

第二层：再看逻辑邻居如何变成待连接任务

这一层看 generic.cc 和 transport.cc 的第一段。

要回答的问题是：

1

ring.prev / ring.next 这些逻辑邻居，怎么告诉 transport 层去连接？

关注函数：

1
2
3
4
5
6

src/transport/generic.cc
  ncclTransportRingConnect
  ncclTransportTreeConnect

src/transport.cc
  ncclTransportP2pConnect

这一层的阶段边界：

1

ncclTransportP2pConnect 不建链。

它只是把待连接的 channel 写进 bitmask：

1
2

comm->connectRecv[peer]
comm->connectSend[peer]

因此它可以视为待连接记录的生成步骤。

第三层：看连接建立函数

这一层进入连接建立逻辑。

要回答的问题是：

1

待连接 bitmask 如何变成 device kernel 能用的 ncclConnInfo？

关注函数：

1
2
3

src/transport.cc
  ncclTransportP2pSetup
  selectTransport

读 ncclTransportP2pSetup 时，可以先将分析范围限定为：

1
2
3
4

rank A 和 rank B
一个 channel
一个 send connector
一个 recv connector

然后按这个顺序看：

1
2
3
4
5
6
7

读 connectSend / connectRecv bitmask
  -> selectTransport
  -> transport.setup
  -> bootstrapSend / bootstrapRecv 交换 ncclConnect
  -> transport.connect
  -> 填 connector->conn，也就是 ncclConnInfo
  -> cudaMemcpyAsync 到 devPeers

这里最重要的是区分两个结构：

1
2
3
4
5
6
7

ncclConnect:
  host 之间交换的临时 handle。
  它携带对端建立连接所需的信息。

ncclConnInfo:
  device kernel 最终使用的连接信息。
  里面有 buffs、head、tail、connFifo、stepSize 等字段。

这两个结构对应不同阶段：ncclConnect 属于 bootstrap 交换阶段，ncclConnInfo 属于 device 侧执行阶段。

第四层：最后再分 transport 看细节

transport 细节可以作为最后一层阅读，第一阶段只关注 setup 和 connect。

阅读顺序：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17

src/transport/p2p.cc
  p2pSendSetup
  p2pRecvSetup
  p2pSendConnect
  p2pRecvConnect

src/transport/shm.cc
  shmSendSetup
  shmRecvSetup
  shmSendConnect
  shmRecvConnect

src/transport/net.cc
  sendSetup
  recvSetup
  sendConnect
  recvConnect

可以先读 P2P，再读 SHM 和 NET：

1
2
3
4
5
6
7
8
9
10
11

P2P:
  GPU 之间可直接或间接访问。
  可能走 direct pointer、CUDA IPC、cuMem，也可能选 intermediate rank。

SHM:
  同 host、共享 /dev/shm 时使用。
  主要是共享内存描述符的创建和 import。

NET:
  跨节点通信的主要路径。
  会涉及 NIC 选择、GDR、PXN、proxy、异步 connect。

proxyProgress 主要用于运行时数据传输的进度推进，可以不纳入初始化建链的第一阶段阅读。

最小 loop

最小阅读 loop 可以只覆盖 ring：

1
2
3
4
5
6

connectRings
  -> ncclTransportRingConnect
  -> ncclTransportP2pConnect
  -> ncclTransportP2pSetup
  -> p2pSendSetup / p2pRecvSetup
  -> p2pSendConnect / p2pRecvConnect

在这个 loop 中暂不展开 tree、CollNet、NVLS、runtime connect。ring 路径建立后，其它路径主要差异为：

1
2
3

邻居集合不同
connIndex 不同
transport 细节不同

对应关系如下：

1
2
3
4
5
6
7

graph 层：生成 channel 的逻辑通信关系
P2pConnect：将逻辑通信关系记录到待连接 bitmask
P2pSetup：处理待连接 bitmask
transport.setup：创建本地资源并生成 ncclConnect
bootstrap：交换 ncclConnect
transport.connect：根据对端 ncclConnect 建立连接
ncclConnInfo：保存 device kernel 使用的连接信息

建链流程

建链流程可以按阶段阅读：先生成 channel 的逻辑邻居，再把逻辑邻居登记成待连接任务，随后选择 transport、交换 handle、填充 ncclConnInfo。

生成逻辑邻居

初始化时，NCCL 先探测机器拓扑，计算 GPU、NIC、CPU 之间路径，然后搜索不同 collective algorithm 对应的 graph。

典型流程在 initTransportsRank 中：

1
2
3
4
5
6
7

ncclTopoGetSystem
ncclTopoComputePaths
ncclTopoTrimSystem
ncclTopoSearchInit
ncclTopoCompute(ringGraph)
ncclTopoCompute(treeGraph)
ncclTopoCompute(collNetGraph / nvlsGraph)

此时 graph 中包含拓扑搜索结果。之后进入：

1
2
3

ncclTopoPreset
bootstrapAllGather(allGather3Data)
ncclTopoPostset

ncclTopoPreset 先基于本节点 intra 排列生成本 rank 视角下的信息，例如：

• topoRanks->ringRecv[c]
• topoRanks->ringSend[c]
• topoRanks->ringPrev[c]
• topoRanks->ringNext[c]
• topoRanks->treeToParent[c]
• topoRanks->treeToChild0[c]
• topoRanks->treeToChild1[c]

然后所有 rank 做一次 bootstrapAllGather，每个 rank 都拿到全局的 topoRanks。

ncclTopoPostset 会调用：

1
2

connectRings
connectTrees

这里会把 channel 中的逻辑邻居写好：

1
2
3
4

channel->ring.prev
channel->ring.next
channel->tree.up
channel->tree.down[]

注意，这一步只回答“谁和谁连”，还没有真的建立 CUDA IPC、SHM 或 NET 连接。

记录待连接 bitmask

以 ring 为例：

1

ncclTransportP2pConnect(comm, c, 1, &channel->ring.prev, 1, &channel->ring.next, 0)

tree 类似，只是 recv/send peer 来自 tree.down[] 和 tree.up。

ncclTransportP2pConnect 不创建 transport 资源，也不交换 handle。它根据 channel id 在两个 bitmask 上记录待连接关系：

1
2

comm->connectRecv[peer] |= 1ULL << channelId
comm->connectSend[peer] |= 1ULL << channelId

因此，ncclTransportP2pConnect 的输出是待连接 bitmask，而不是已建立的连接。

连接建立发生在随后调用的：

1

ncclTransportP2pSetup(comm, graph, connIndex)

建立连接并写入 ncclConnInfo

ncclTransportP2pSetup 可以拆成一个循环：

更具体一点：

1. 按 rank 距离枚举 peer。
2. 读取 connectRecv[recvPeer] 和 connectSend[sendPeer]，知道哪些 channel 需要连接。
3. 对每条待连接边调用 selectTransport。
4. selectTransport 按顺序尝试 P2P、SHM、NET、COLLNET。
5. 命中某个 transport 后调用它的 setup。
6. setup 创建本地资源，并把远端需要的信息写到 ncclConnect。
7. 通过 bootstrap 把 ncclConnect 交换给对端。
8. 调用对应 transport 的 connect，用对端 handle 去 import/map 资源。
9. 填好 connector->conn，也就是 ncclConnInfo。
10. 把 ncclConnInfo 拷到 device side 的 devPeers。
11. 最后做一次 bootstrap 同步，清理 connectSend/connectRecv bitmask。

ncclConnInfo 是 device kernel 使用的信息，核心字段包括：

• buffs[]
• head
• tail
• connFifo
• stepSize
• flags
• 部分 NET/GDR 场景下的 net device handle 和 memory handle

选择 transport

transport 抽象在 src/include/transport.h：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

struct ncclTransportComm {
  ncclResult_t (*setup)(...);
  ncclResult_t (*connect)(...);
  ncclResult_t (*free)(...);
  ncclResult_t (*proxySharedInit)(...);
  ncclResult_t (*proxySetup)(...);
  ncclResult_t (*proxyConnect)(...);
  ncclResult_t (*proxyFree)(...);
  ncclResult_t (*proxyProgress)(...);
  ncclResult_t (*proxyRegister)(...);
  ncclResult_t (*proxyDeregister)(...);
};

struct ncclTransport {
  const char name[8];
  ncclResult_t (*canConnect)(...);
  struct ncclTransportComm send;
  struct ncclTransportComm recv;
};

全局 transport 顺序在 src/transport.cc：

1
2
3
4

struct ncclTransport* ncclTransports[NTRANSPORTS + 1] = {
  &p2pTransport, &shmTransport, &netTransport, &collNetTransport,
  &profilerTransport
};

selectTransport 会按这个顺序调用 canConnect，第一个返回可用的 transport 就会被选中。

因此普通点对点边的选择顺序为：

1
2
3

优先 P2P
其次 SHM
最后 NET

P2P transport

src/transport/p2p.cc 处理 GPU 之间可以直接或间接访问的场景。

setup 阶段：

• 判断是否可以 CUDA P2P。
• 判断使用 read 还是 write。
• 判断是否需要 intermediate rank。
• 同进程可能用 direct pointer。
• 跨进程可能用 CUDA IPC 或 cuMem handle。
• 通过本地 proxy 分配或导出可共享 buffer。
• 把 p2pConnectInfo 写入 ncclConnect。

connect 阶段：

• import/map 对端 buffer。
• 设置本端 conn.buffs[]。
• 设置 head/tail。
• 设置 ptrExchange/redOpArgExchange 等辅助字段。
• 设置 proxyProgress。

P2P 连接不一定是 A 与 B 直接连接。p2pGetInfo 可能选择 intermediate rank，形成 indirect P2P path。

SHM transport

src/transport/shm.cc 处理同 host 且共享 /dev/shm 的场景。

shmCanConnect 主要检查：

• NCCL_SHM_DISABLE 是否关闭 SHM。
• topo 是否要求走 NET。
• 两个 rank 是否在同 host。
• 两个 rank 是否有共同的 shm device。

setup 阶段通过 SHM proxy 创建共享内存，把描述符写入 shmConnectInfo。

connect 阶段 import 对端共享内存，然后设置：

1
2
3
4

conn.buffs[]
conn.head
conn.tail
conn.stepSize

NET transport

src/transport/net.cc 是跨节点通信的主要 transport，也可能在同节点但拓扑要求走 NET 时使用。

setup 阶段会：

• 通过 ncclTopoGetNetDev 选择 NIC。
• 判断是否启用 GDR。
• 判断是否需要 PXN proxy。
• 通过 ncclProxyConnect 连到对应 proxy rank。
• 对 proxy 发送 ncclProxyMsgSetup。
• 把 proxy rank 和 GDR 信息放进 ncclConnect。

NET connect 阶段可能异步进行：

1
2

ncclProxyCallAsync(ncclProxyMsgConnect)
ncclPollProxyResponse

如果还没完成，会返回 ncclInProgress。外层 ncclTransportP2pSetup 因此会循环 polling，直到所有 channel 都完成连接。

NET connect 完成后，会 map host/device/shared memory，并设置：

1
2
3
4
5
6

conn.head
conn.tail
conn.connFifo
conn.buffs[]
conn.netDeviceHandle
conn.mhandles[]

runtime connect

另一个分支是 runtime connect：

1

comm->runtimeConn = comm->cuMemSupport && ncclParamRuntimeConnect();

如果 runtime connect 开启，初始化阶段不会把 ring/tree 都预连接好，而是在 group/enqueue 阶段按需连接。

相关入口：

• ncclP2PPreconnectFunc
• ncclCollPreconnect

这也是为什么有时候读 initTransportsRank 会发现某些连接没有在 init 阶段直接发生。

通信流程

建链完成后，NCCL 不会在每次通信时重新选择 transport。建链阶段已经把每个 channel、peer、connIndex 对应的连接信息写入 connector->conn，并同步到 device side 的 devPeers。通信阶段复用这些连接信息，主要完成三件事：

1
2
3

host 侧生成本次通信的 work
device kernel 按算法执行 copy / reduce / send / recv
proxy 线程在需要时推进 NET 或其它需要 CPU 参与的 transport

下图展示一次通信对建链产物的使用方式：

Host 侧生成 work

一次 collective 或 P2P API 调用先被放入 communicator 的 planner。group launch 阶段会调用：

1
2
3

ncclLaunchPrepare
  -> scheduleCollTasksToPlan / scheduleP2pTasksToPlan
  -> finishPlan

ncclLaunchPrepare 的输出是 ncclKernelPlan。plan 中包含：

• 本次通信使用哪些 channel。
• 每个 channel 上的 device work。
• kernel 参数和 work buffer。
• 需要 proxy 推进时的 proxy op。

随后 launch 阶段会执行：

1
2
3
4
5
6
7
8

uploadWork
  -> 把 ncclDevWork 写入 kernel args 或 work FIFO

uploadProxyOps / ncclProxyStart
  -> 如果 transport 需要 proxy，则把 proxy op 投递给 proxy progress thread

ncclLaunchKernel
  -> 按 channel 数启动 NCCL device kernel

在这个阶段，host 侧主要负责“把本次通信描述清楚”。数据搬运发生在 device kernel 和 proxy progress 中。

Device kernel 消费连接信息

device kernel 拿到 work 后，会根据 collective 类型、algorithm、protocol 进入对应实现。

以 ring allreduce 为例，核心路径在 src/device/all_reduce.h 的 runRing。它使用建链阶段写好的：

1
2

ncclShmem.channel.ring.prev
ncclShmem.channel.ring.next

并构造 Primitives。ring allreduce 的执行可以分成两个阶段：

1
2
3
4
5
6
7
8

reduce-scatter:
  directSend
  recvReduceDirectSend
  recvReduceCopyDirectSend

all-gather:
  directRecvCopyDirectSend
  directRecv

这些 primitive 最终会进入 src/device/prims_simple.h 中的 genericOp。genericOp 的基本循环是：

1
2
3
4
5
6
7
8
9
10

waitPeer
  -> 等待对端 step 状态满足条件
  -> 根据 connFifo / directBuff / buffs 计算本 slice 使用的地址

reduceCopy
  -> 执行 copy、reduce 或 reduce + copy

postPeer
  -> 更新 step counter
  -> 通知对端当前 slice 已经生产或消费完成

对于用户显式的 ncclSend / ncclRecv，device 侧路径在 src/device/sendrecv.h。send 侧循环调用 directSend，recv 侧循环调用 directRecv，同样通过 Primitives 使用建链阶段准备好的连接。

step 和 FIFO

ncclConnInfo 中的 buffs、head、tail、connFifo、stepSize 是通信阶段的关键状态。

可以把它理解为一个带 credit 的环形 FIFO：

1
2
3
4
5
6

NCCL_STEPS 个 slot 循环使用
每个 slice 对应一个或多个 step
发送方等待可写 slot
接收方等待可读 slot
数据写入或读出后更新 step counter
head / tail / connFifo 用于表达空间、数据大小和可见性

因此，通信阶段不是简单地调用一次 memcpy。它是按照 slice/chunk 粒度，在 channel 上反复执行：

1
2
3
4
5

等待可用 step
确定 buffer 地址
copy / reduce
发布 step
进入下一个 slice

Proxy 推进 NET 等传输

如果当前连接是 GPU 可以直接访问的 P2P 路径，device primitive 可以直接读写对端可见的 buffer。

如果当前连接需要 CPU proxy，例如 NET transport，device kernel 和 proxy thread 会共同推进传输：

1
2
3
4
5
6
7
8
9
10
11
12
13

device kernel:
  写入 staging buffer
  更新 tail 或 connFifo

proxy progress thread:
  发现 GPU 数据 ready
  调用 ncclNet->isend / ncclNet->irecv
  轮询 ncclNet->test
  网络完成后更新 head / tail / connFifo

device kernel:
  观察 step 状态变化
  继续消费下一个 slice

对应代码路径：

1
2
3
4
5
6
7
8

src/proxy.cc
  ncclProxyStart
  ncclProxyProgress
  progressOps

src/transport/net.cc
  sendProxyProgress
  recvProxyProgress

sendProxyProgress 主要做三件事：

1
2
3

给 GPU 投递可用 buffer
等待 GPU 写完并设置 connFifo size / tail
调用 ncclNet->isend，完成后释放 FIFO slot

recvProxyProgress 主要做三件事：

1
2
3

把接收 buffer 投递给 ncclNet->irecv
等待网络接收完成
更新 GPU 可见的状态，让 device kernel 继续读取

建链和通信的边界

二者的边界可以概括为：

1
2
3
4
5
6

建链阶段：
  选择 transport，交换 handle，填 ncclConnInfo。

通信阶段：
  复用 ncclConnInfo，把本次 API 调用切成 channel work，
  由 device kernel 和 proxy progress 按 step/FIFO 协议完成数据传输。

小结

NCCL 的建链和通信可以按两个阶段理解。

建链阶段可以概括为：

1
2
3
4

topology 决定 channel 的逻辑邻居，
P2pConnect 把待连接边登记到 bitmask，
P2pSetup 选择 transport、交换 handle、填 ncclConnInfo，
最后把 conn 信息拷到 device，让 kernel 直接通信。

通信阶段可以概括为：

1
2
3
4

host 把 API 调用切成 ncclKernelPlan，
device kernel 按 channel 执行 collective 或 P2P work，
Primitives 通过 waitPeer / reduceCopy / postPeer 消费 ncclConnInfo，
NET 等路径由 proxy progress 线程协同推进。

源码结构对应关系：

• graph/connect.cc 负责生成 channel 的逻辑通信关系。
• transport.cc 回答“这些边什么时候建、怎么调度建”。
• transport/p2p.cc、transport/shm.cc、transport/net.cc 回答“具体 transport 如何创建资源、交换 handle、填 device 可用的连接信息”。
• enqueue.cc 和 group.cc 负责把一次 API 调用变成 kernel plan。
• device/all_reduce.h、device/sendrecv.h 和 device/prims_simple.h 负责 device 侧 copy、reduce 和 step/FIFO 协议。
• proxy.cc 和 transport/net.cc 负责需要 proxy 参与的运行时传输。

按上述边界划分后，NCCL 可以拆成两条相互衔接的主线：建链阶段准备连接，通信阶段复用连接。

MAI-Thinking-1 的技术报告标题是 MAI-Thinking-1: Building a Hill-Climbing Machine。

官方还有一个介绍页：Introducing MAI-Thinking-1，以及 MAI-Thinking-1 Model Card。

一句话概括：这篇报告围绕 Microsoft AI 所说的 hill-climbing machine 展开，内容覆盖数据、MoE 架构、训练系统、RL recipe、agent 环境、评测和安全红队，重点是模型能力迭代体系，而不只是单个模型的 benchmark 结果。

MAI-Thinking-1 本身是一个 35B active / 1T total parameter 的 sparse MoE reasoning model。报告强调它是 from scratch 训练的，没有用第三方模型蒸馏；pre-training 使用 30T tokens，数据来自公开和授权的人类生成内容，并且在 pre-training 阶段不使用语言模型生成的 synthetic data。

需要区分的是：报告里说的 “不蒸馏” 主要指不从第三方模型继承能力；RL 阶段内部仍然使用 self-distillation 来恢复或延续 RL climb。

本文主要整理几个点：

• 模型结构：MAI-Base-1 的 MoE、attention、LatentMoE 和 dropless MoE；
• 训练基础设施：YOLO 训练框架、硬件集群、调度和运维；
• Pre-training：scaling ladder、数据 composition 和训练阶段；
• Post-training / RL：RL climb、STEM / agentic / helpfulness-safety 数据与训练；
• Evaluation：benchmark、人类偏好评测和安全评估。

模型结构

MAI-Thinking-1 的 base model 叫 MAI-Base-1，是 decoder-only Transformer + sparse MoE。

几个核心配置：

这个结构主要服务于训练效率、推理效率和大规模训练稳定性。

Attention 部分没有用每层 full attention，而是 5 层 local attention 搭配 1 层 global attention。local attention 用 RoPE，window size 是 512；global attention 不用 position encoding。这样可以降低训练 attention 计算量，也能减少推理时 KV cache 的规模。

FFN 部分采用 dense FFN 和 MoE 交替。它不是 every-layer MoE，而是把高稀疏 MoE 和零稀疏 dense layer 配在一起。报告里的说法是，这种 interleaved layout 在 scaling 上接近更均匀的 medium-sparsity MoE，但 wall-clock 训练效率更好。

MoE 使用 LatentMoE：先共享 down-projection，再做 all-to-all dispatch；routing 仍然基于原始 representation，每个 compressed representation 被路由到 512 个专家中的 8 个。报告将其作为降低专家通信和计算成本的一项设计。

另一个工程细节是 dropless MoE。很多 MoE 实现会设置 expert capacity，超出 capacity 的 token 被 drop。报告里说他们最终收敛到 fully dropless MoE，并支持 variable message size all-to-all 和 bounded memory usage。这个设计与训练稳定性有关：如果 token dropping 存在，routing / load balance 的结果会受到 capacity 设置影响。

训练基础设施

训练基础设施部分主要包括 YOLO 训练框架、底层硬件集群、Kubernetes/Kueue/Ray 调度运行时，以及围绕 goodput 的稳定性和可观测性体系。

YOLO 训练框架

报告里的训练系统叫 YOLO，展开是 You Only Launch Once。

YOLO 是 Microsoft AI 自研的大规模训练框架，基于 PyTorch，覆盖 pre-training、mid-training、SFT 和 RL training。它实现了 model definition、sharding、optimizer、dataloader、checkpointing 等核心训练循环。

报告没有说明 YOLO 与 DeepSpeed 的实现关系；虽然二者都覆盖 ZeRO、MoE、sharding、checkpointing 等大规模训练问题域，但不能据此判断 YOLO 是否复用了 DeepSpeed。

报告列出的训练系统能力包括：

• 自定义 FP8 GEMM、Grouped GEMM、quantization kernels；
• MoE 支持 dropless / capacity-capped、多种 load-balancing 策略、router replay；
• expert dispatch / compute / collect 做 pipeline overlap；
• activation checkpointing + activation offloading；
• bitwise reproducibility；
• checkpoint 保存 model weights、optimizer state、FP8 scaling history、dataloader progress、RNG；
• 把 goodput 当成生产 KPI，而不仅仅是 MFU。

报告里提到 MAI-Base-1 pre-training 在 8K GPUs 上达到 90.0% goodput。goodput 不同于单纯的 MFU，它会受到 crashloop、node failure、link flap、OOM、checkpoint stall、recompute、slow startup、调度延迟、恢复后 MFU 下降等因素影响。

从 infra 角度看，这部分说明报告关注的不只是单步训练吞吐，也包括故障恢复、可复现性、checkpoint 和长周期训练的有效利用率。

报告给出的 goodput 定义是：

1

goodput = ideal training duration / actual wall-clock duration

其中 ideal training duration 可以理解为：如果训练过程一直以目标状态稳定 stepping，没有失败、重算、启动等待、checkpoint stall、MFU drop 等 overhead，完成同样训练进度所需的理想时间。actual wall-clock duration 则是实际从训练开始到完成所花的墙钟时间。

也可以按 overhead 分解成：

1
2
3

actual wall-clock duration = ideal training duration + total overhead

goodput = ideal training duration / (ideal training duration + total overhead)

因此：

1

overhead ratio = 1 - goodput

MAI-Base-1 的 90.0% goodput 表示：实际消耗的 wall-clock time 中，大约 90% 对应理想训练进度，约 10% 是各种 overhead。报告进一步把 overhead 拆成 recomputation、non-stepping time、MFU drop overhead 等类别。

这个指标和 MFU 的关系是：MFU 衡量 stepping 时 GPU 理论算力被模型计算利用了多少；goodput 衡量端到端训练墙钟时间里有多少真正转化成有效训练进度。因此，一个 run 可以有较高 MFU，但如果频繁失败、重启、重算或 checkpoint 卡住，goodput 仍然会很低。

报告还披露了几组训练系统指标：

• MAI-Base-1 pre-training 在 8K GPUs 上达到 90.0% goodput；
• total overhead 降到 51 hours；
• recomputation 为 6.5 hours，占 overhead 的 15%；
• non-stepping time 为 14 hours，占 overhead 的 27%；
• MFU drop overhead 为 18 hours，占 overhead 的 35%，是最终 run 中最大的剩余 overhead 类别；
• 在架构演进中，早期 GB200 NVL64 设置的 MFU 从 18% 提升到 22%；
• GPU Direct RDMA 约带来 1.1x end-to-end training time 改善；
• 自定义 block-sparse attention backend 约带来 1.06x step time 改善；
• ZeRO-2 相关优化约带来 1.03x end-to-end 改善；
• 自定义 Triton expert encode kernel 将 HBM utilization 从约 10% 提升到约 80%，并带来约 1.03x end-to-end 改善。

推理部署侧，报告称 MAI-Thinking-1 在 MAIA-200 上实现后，相比 GB200-based deployment，在相同 rack power budget 下 token generation throughput 高 40% 以上。这个指标是 performance per watt 口径，不是单卡峰值吞吐。

硬件和集群架构

MAI-Thinking-1 的主训练硬件是 Microsoft Azure 平台内 Microsoft-operated 的单站点 GB200 集群。这个部分的信息主要来自报告正文的 Section 6 Cluster Environment 和 Appendix K Cluster Environment Details；外部链接只用于补充硬件、网络和调度组件背景。

报告没有说明使用了 Azure Machine Learning。它披露的是：MAI-Base-1 在 Azure platform 内的 Microsoft-operated cluster 上，从 8K GB200 GPUs 开始做 from-scratch pre-training；调度和运行时栈包括 Kubernetes、Kueue、MAI control plane、Ray、NCCL 和 YOLO 训练框架。下文按 Azure-based first-party infrastructure 理解。

下图是基于报告文字整理的系统架构示意。

图中 GHR 是 guest health report，指节点内部运行环境上报的健康信息，用于辅助判断节点是否需要进入 vendor maintenance。NPD 是 Kubernetes Node Problem Detector，负责把节点侧异常上报为 Kubernetes node conditions 或 events。admitted job 指已经通过集群准入 / 调度流程，并获得相应计算资源的训练作业。

报告里的硬件口径如下：

GB200 和 GB300 集群部署在 Microsoft first-party datacenters。这里的 first-party datacenters 可理解为 Microsoft 自有或直接运营的 Azure 数据中心资源。

报告还提到，这些集群通过 Azure 团队共同维护的 custom images 暴露给 MAI。从 Appendix K 的 cluster provisioning、node lifecycle、certification 和 telemetry 语境看，这里的 custom images 更像是节点级操作系统镜像，可能包含 OS、GPU driver、RDMA/NCCL 相关组件、诊断与 telemetry agent 等基础软件栈；报告没有明确说明它是 OS image 还是 container image。

主训练被放在一个 single logical cluster、one site 上，主要是为了降低实验方差：同一代 accelerator、稳定 rack health、稳定 scheduler 行为、可预测 storage path。

硬件拓扑上，GB200 / GB300 系统以 rack-scale NVL72 为单位部署：

• 每个 rack 是一个 72-GPU NVLink domain；
• NVLink / NVSwitch 负责 rack 内 scale-up 高带宽通信；
• rack 间 scale-out 通信用 InfiniBand RDMA；
• 为了训练稳定性，报告中实际使用 64 GPUs per rack，即 NVL64，保留 spare capacity 来容忍 node failure 和 unhealthy devices；
• H100 系统仍在 lab 环境里使用，形态是 8-GPU nodes，node-local NVLink/NVSwitch，跨节点 InfiniBand。

可以简化理解为：

1
2
3
4
5
6
7
8
9
10

GB200 rack / NVL72
  72 GPUs in one NVLink domain
  report training placement uses 64 GPUs per rack (NVL64)

multiple racks
  connected by InfiniBand RDMA

large training job
  keep expert all-to-all inside NVL64
  use cross-rack InfiniBand mainly for data parallel communication

这个拓扑也影响了模型并行策略。为了提高 GEMM efficiency，MAI-Base-1 选择：

• expert parallelism，EP = 64；
• tensor parallelism，TP = 1；
• expert all-to-all communication 保持在 NVL64 domain 内；
• cross-rack InfiniBand 用于 data parallel communication，比如 parameter all-gather 和 gradient reduce-scatter；
• pre-training 和 mid-training 1 使用 EP=64 + ZeRO-2；
• mid-training 2 启用 ZeRO-3 / FSDP；
• mid-training 阶段使用 context parallelism。

集群切分和逻辑集群

Appendix K 里有一句关键描述：每个 site 会被切分成多个 Kubernetes clusters，通常是一栋 datacenter building 对应一个 Kubernetes cluster。

可以理解为：

1
2
3
4

physical site / datacenter campus
  datacenter building A -> Kubernetes cluster A
  datacenter building B -> Kubernetes cluster B
  datacenter building C -> Kubernetes cluster C

这种切分不是说一个训练任务只能在一栋楼里运行，而是把物理基础设施按 building 边界组织成多个 Kubernetes 管理单元。报告同时提到 large jobs 在需要时可以跨越单个 Kubernetes cluster 的边界：nodes are universally routable across the compute environment，workload pods use host networking，以减少 overlay network overhead。

这里的几个层次可以分开看：

报告称主训练放在一个 single logical cluster、one site 上，目标是降低实验方差：同一代 accelerator、稳定 rack health、稳定 scheduler 行为、可预测 storage path。这里的 logical cluster 不应直接等同于单个 Kubernetes cluster；Appendix K 同时说明一个 site 通常会按 datacenter building 切分为多个 Kubernetes clusters。因此，8K GB200 是否全部位于同一个 Kubernetes cluster 内，报告没有明确披露。

调度和控制面

控制面和调度架构可以拆成：

• Kubernetes 维护 cluster state；
• Kueue 负责 quota、admission、priority、preemption 和 topology-aware placement；
• MAI cluster-local control plane 管理 reservation、rack topology、quota coherence 和 scheduling-readiness gates；
• Ray 在 admitted jobs 内执行 distributed runtime；
• MAI drivers 把调度得到的 topology 转成 actor placement、communication groups 和 NCCL clique configuration。

这里的 MAI control plane 不是报告中披露的公开 Kubernetes 插件名。按 Appendix K.3 的描述，它更像一组 MAI 内部的 cluster-local controllers：不替代 scheduler，而是维护 scheduler 所需的 reservation、rack topology、quota coherence、scheduling-readiness gates 等状态。是否实现为 Kubernetes CRD/controller、scheduler plugin 或其他内部组件，报告没有进一步披露。

Kueue 负责 admission 和 topology-aware placement。MAI cluster-local control plane 则维护 Kueue / scheduler 做决策所需的状态，例如 rack reservation、topology labels、scheduling-readiness gates。报告特别提到 rack fragmentation 风险：许多小任务如果随机占满不同 rack，会让后续大任务很难拿到连续、拓扑紧凑的容量。为此，cluster-local control plane 会维护 soft rack reservations。队列可以有 preferred racks；空闲时可以借用容量，需要时再通过 reclaimWithinCohort reclaim reserved racks。

Ray runtime 和训练作业

Ray 是 admitted job 内部的 distributed runtime。这里的 admitted job 指已经通过集群准入 / 调度流程，并获得相应计算资源的训练作业。Kueue 完成 admission 和 placement 后，MAI drivers 会把得到的 topology 转成 Ray actor placement、communication groups 和 NCCL clique configuration。

不同类型任务的 actor 需求不同：

• pre-training jobs 主要要求 strict learner availability；
• RL jobs 会包含 learners、inference servers、rollout workers、routers 等多种 actor；
• MAI drivers 负责监控 actor liveness、协调训练循环，并维护异步组件之间的 checkpoint consistency。

这部分和 RL training 尤其相关。RL 不是单一同步训练循环，而是 learner、推理服务、rollout、reward / grader 等多类组件共同工作。报告把 Ray 放在 admitted jobs 内部；按该描述，Ray 的作用范围主要在作业内，用于 actor 编排和运行时管理，集群级排队、准入和资源分配仍由 Kubernetes、Kueue 以及 MAI control plane 等组件承担。

Certification 和节点生命周期

报告强调 physical topology 和 hardware health 是 first-class scheduling state。节点不是 provisioned 就可用，而是要经过 certification。certification 的目的，是防止坏节点、退化链路、边缘状态存储和 silent-corruption 风险进入生产训练池。

certification 分层进行：

节点生命周期见报告 Figure 26：

状态迁移可以直接理解为：

• 新节点或修复后的节点先进入 Init，再由 certification controller 进入 Certifying；
• certification 通过后进入 Available，失败则进入 Impaired；
• Available 节点在出现 NPD condition 或 manual drain 后会进入 Impaired；
• 如果是误报，可以从 Impaired 回到 Available；
• 如果需要自动修复，则从 Impaired 进入 Auto Remediating，例如 auto reboot、reset 或 soft drain，之后回到 Init 重新认证；
• 如果自动修复不足以解决问题，则通过 GHR 进入 Vendor Maintenance；
• vendor 维修成功后进入 Repaired，再回到 Init；维修失败则进入 Decommissioned。

这里 Impaired 表示节点已被判定为不健康或不适合继续进入生产训练池；Auto Remediating 表示系统先尝试自动修复，例如 reboot、soft drain 或 reset。NPD 是 Kubernetes Node Problem Detector，负责把节点侧异常上报为 Kubernetes node conditions 或 events。GHR 可理解为 guest health report，即节点运行环境内部上报的健康信息，通常来自 guest OS 或节点级 agent，用于补充平台侧硬件健康信号。Runtime monitoring 会根据 NPD conditions、XID errors、ECC thresholds、NVLink degradation、InfiniBand link flaps、storage faults 等信号触发状态迁移。修复后的节点不会直接回到 Available，而是回到 Init 并重新经过 certification。

异常处理时序可以整理成：

sequenceDiagram
participant Telemetry as Telemetry / Runtime Monitoring
participant Controller as Certification / Remediation Controller
participant Scheduler as Scheduler / Kueue
participant Node as Node
participant Vendor as Vendor / Datacenter Maintenance

Telemetry->>Controller: XID / ECC / NVLink / IB / storage fault
Controller->>Scheduler: mark node unschedulable / drain
Controller->>Node: attempt auto remediation

alt transient issue fixed
  Node-->>Controller: reboot / reset / soft drain succeeds
  Controller->>Node: reset to Init
  Controller->>Node: run certification
  Controller->>Scheduler: mark node Available
else persistent hardware issue
  Controller->>Vendor: GHR / guest health report
  Vendor->>Node: repair or replace
  Node-->>Controller: repaired
  Controller->>Node: reset to Init
  Controller->>Node: re-run certification
  Controller->>Scheduler: mark node Available if cert-pass
else maintenance fail
  Vendor->>Controller: maintenance fail
  Controller->>Scheduler: decommission node
end

Telemetry 和 observability

报告把 observability 放进控制回路，而不是只做 dashboard。硬件 telemetry、fabric health、storage behavior、scheduling state 和 job progress 会决定容量是否 admitted、drained、remediated 或 returned to service。

硬件健康信号包括：

• GPU XID；
• ECC；
• thermals、power、clock throttling；
• NVLink state、NVLink bit-error rate、chip-to-chip links；
• InfiniBand device state；
• local NVMe health；
• PCIe errors；
• driver state。

这些信号会转成 Kubernetes node conditions，再进入 scheduling、triage、drain 和 remediation controllers。

作业可观测性则横跨 Kueue、Kubernetes、Ray、training logs 和 experiment metadata。报告提到 operator 可以按 namespace、pod、job、restart index 查看 queue、priority、admission state、node placement、worker readiness、restart count、training configuration、step progress 和 scoped logs。这样可以区分 scheduling delay、runtime failure、node failure、storage degradation 和 application-level stalls。

Telemetry 存储和查询也分层：

因此，硬件架构不只是 “8K GB200 GPUs”，而是一套围绕 usable training capacity 设计的系统：rack 内 NVLink 负责高带宽局部通信，rack 间 InfiniBand 负责扩展到多 rack；调度器尽量保持 locality；certification 和 telemetry 防止坏节点、坏链路和 silent corruption 进入训练池。

参考信息

本节事实来源主要是 MAI-Thinking-1 技术报告的 Section 6 Cluster Environment 和 Appendix K Cluster Environment Details。下面这些外部链接只作为背景参考，用来解释报告中出现的硬件、网络和调度组件，不作为 MAI 训练细节的独立来源。

Pre-training

Pre-training 部分主要整理 scaling ladder、pre-training 数据 composition，以及 pre-training / mid-training 的训练阶段。

Scaling ladder

报告反复强调 scaling ladder：架构和数据决策不只看单个小规模实验，而要看收益能否沿着 scale 稳定成立。

他们用 scaling ladder 做架构和数据消融：对不同 model size，用固定的 tokens per active parameter 训练，比较 scaling curve。多数 architecture ablation 在接近 Chinchilla optimal 的 100-200 TPP 做，而主训练会 over-train 到 500-1000 TPP，让模型更适合高频推理场景。

这个方法的前提是：小模型上的改进不一定能迁移到大模型；某个数据 mixture 在小规模上更好，也不代表在大规模上排序不变。因此，报告把可扩展验证放在 pre-training 决策的中心。

这也是 “hill-climbing machine” 的组成部分：架构、数据和训练系统都通过 ladder + efficiency gain 做评估，以支持持续迭代。

Pre-training 数据

MAI-Base-1 使用 30T tokens 预训练。数据来源包括：

• web HTML；
• web PDFs；
• public GitHub code；
• books and journals；
• academic papers；
• news；
• multilingual text；
• domain-specific materials。

报告披露了几个数据治理口径：

• 不使用 open source training datasets；
• 不使用语言模型生成的 synthetic data 做 pre-training；
• 尽力移除采集源里的 AI-generated content；
• 排除常见机器学习数据站点和仓库，比如 huggingface.co 一类来源；
• 不使用 Microsoft 产品和服务里的 private customer data，除非用户明确 opt in 或适用协议允许；
• 对整个 corpus 做 PII-risk 和 safety filtering。

报告披露的知识截止日期也比较细：

从报告披露看，MAI-Base-1 的 pre-training 数据策略强调 clean、licensed 和 human-generated。

预训练数据 composition 的数字如下：

这里有几个指标比较关键：

• code 占 54.6%，是最大的数据来源；
• STEM + Math 合计 21.2%，并且 Math 的平均采样 epoch 最高，达到 5.28x；
• Web text 和 PDFs 的可用 unique tokens 没有被完整耗尽，平均 epoch 分别是 0.55x 和 0.53x；
• multilingual other 只有 1.6% training mix，但报告说明 domain-specific multilingual data 会被计入其他类别。

mid-training 的数据仍来自 pre-training corpus，不引入新的 synthetic source。报告披露的目标 mixture 是：code 55%，STEM/math 35%，background sources 10%。

训练阶段的规格如下：

也就是说，MAI-Base-1 先在 16K context 上完成 30T tokens 的主预训练，再通过 3.55T tokens 的 mid-training 把 context 扩展到 256K。

Post-training / RL

Post-training / RL 部分包括通用 RL recipe、STEM climb、agentic coding / tool use、helpfulness / safety，以及最后的 consolidation。

RL climb

pre-training 和 mid-training 提供模型预测能力和知识；解题策略、推理 token 分配、工具使用、偏好和安全约束主要在 RL climb 阶段处理。

MAI-Thinking-1 的 RL 从没有 reasoning trace 暴露的 checkpoint 开始，目标是让模型从零发展 reasoning 能力。报告里把稳定 RL 爬坡归因到三个机制：

• 对 GRPO 做两个简单但关键的调整；
• self-distillation，用于 crash 或 base policy 更新后恢复 climb；
• infra 改进，减少 training 和 inference 之间的 numerical mismatch。

他们训练了三个 specialist：

• STEM / competitive code；
• agentic coding / tool use；
• helpfulness / safety。

之后再把这些 specialist consolidation 到单个 MAI-Thinking-1 模型。

GRPO 有两个主要改动。

第一个是 adaptive entropy control。它不是显式加 entropy bonus，而是根据目标 entropy 在线调整 clipping bound。如果 entropy 太低，就放宽上界，让 policy 能更积极地增加 alternative tokens 的概率；如果 entropy 足够高，就收紧 trust region。

第二个是 outer ratio clip。原始 PPO / GRPO 的某些分支在 “往正确方向修正” 时不 clip，但报告发现这些 unclipped branches 会导致 catastrophic gradient-norm spikes，于是加了 hard outer clip。

reward 结构也保持统一：

1

R = R_task + w_lang * R_lang - w_len * R_len

也就是 task reward 加语言一致性奖励，再减长度惩罚。语言一致性奖励用于缓解长上下文 RL 中的语言漂移；长度惩罚用于控制 reasoning 长度。

RL 训练的部分超参数也有披露：

• top-p sampling 使用 p = 0.97；
• 早期最大 rollout length capped at 8K tokens；
• 随训练推进，rollout length 按 2 的幂次扩展，最终到 128K tokens；
• 在 128K extension stage 移除 length penalty，即 w_len = 0；
• problem sampling 设置 G = 128 total rollouts，G_early = 16；
• early pass-rate filtering 使用 [0.05, 0.8]，常规 pass-rate filtering 使用 [0.1, 0.8]；
• 每次 inference model update 之间做 5 个 gradient steps；
• 超过 8 次 inference updates stale 的 rollout 会被丢弃，也就是最多落后 40 gradient steps；
• RL 阶段 global MoE load balancing coefficient 为 1e-5；
• self-distillation SFT 使用 128K sequence length、global batch size 2048、AdamW weight decay 0.001、maximum learning rate 1.7e-5、minimum learning rate 5.2e-6、warmup ratio 2%；
• self-distillation dropout rate 为 0.15，MoE load balancing coefficient 为 1e-2。

STEM 数据

STEM climb 使用 STEM Mix。报告说他们处理了数百万文档，产出超过 5M samples，其中最难部分超过 550k 个 (q, a) pairs。

这个 pipeline 将 textbooks、academic PDFs、forum discussions、competition archives、vendor problems 等异构来源转成可验证的问答对。

流程包括：

• OCR 和 cleanup；
• 去 boilerplate，规范化文本；
• 构建层级结构；
• LLM 标注 question / answer spans；
• 对分离的题目和答案做 QA pairing；
• 标注题型、topic taxonomy、PII、answer leakage；
• 把选择题、证明题等重写成 open-ended；
• 多 solver pass@k 解题；
• consensus grading；
• difficulty rating；
• 丢掉 faulty ground truth。

这个部分的重点是构造可验证、难度合适、topic 多样且 ground truth 可靠的训练信号。报告把 verifier / grader / data pipeline 作为 STEM RL 数据构造的核心环节。

STEM Mix 的样本分布也有披露。

按原始题型：

按学科 taxonomy：

Multiple-choice 和 proof problems 会在 ingestion 过程中转换成 open-ended form；转换不可行的样本会被丢弃，但报告保留了少量 multiple-choice problems，让模型仍然熟悉该格式。

Agentic coding 和 tool use