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

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。

K8s 层：先把物理资源圈对

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 启动：注册 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 分支和特性分支的分叉点。

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

# When topology information is available:
# sort by (domain_min_topo_rank, topo_rank, 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
4
5

# Per-logical-node extra Ray resource requirements.
# Each dict is merged into every bundle spec for that node,
# pinning it to a physical domain.
# Example: [{"nvlink_domain_<uuid>": 0.001}] * 16
# pins 16 nodes to a single 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 尺寸匹配。

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 和重启恢复共同组成一个稳定闭环。

报告里 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 主要用于运行时数据传输的进度推进，可以不纳入初始化建链的第一阶段阅读。

最小闭环

最小阅读闭环可以只覆盖 ring：

1
2
3
4
5
6

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

在这个闭环中暂不展开 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

Agentic climb 训练模型在外部环境里做多步任务：读代码、改文件、跑测试、观察失败、修复方案、再尝试。

报告里的 Sandbox Execution Environment（SEE）会为每个 agentic task 启一个 fresh container，任务完成后销毁。容器默认 network-isolated，以保证可复现，避免 rate limit、外部服务波动等副作用。如果确实需要网络，比如安装包，则通过 caching proxy 和 domain allowlist 管理。

SWE RL problem 会被打包成自包含容器镜像：

• repo checkout 到指定 commit；
• 依赖预装；
• problem statement；
• unit tests / grader；
• Bash tool；
• String replace editor。

模型通过 tool call 和容器交互，结束后 grader 在同一个容器里跑测试并给 reward。

这个设计把 SWE-bench 式任务扩展为可规模化生产的 RL environment。它不是只用于评测的 benchmark，而是用于训练的环境工厂。

tool-use 环境则模拟企业和消费场景中的 API / MCP 交互。每个问题包含 query、tool schemas、initial state 和 grader。有些环境里单个任务可用工具超过 50 个，用来训练模型高效选择合适工具。

他们还会合成 closed-world tool-use environment：生成数据库、tool definition、verifiable task，然后执行和去重。这里 synthetic data 用在 RL / tool-use 环境构造上，而不是 pre-training。

Helpfulness 和 safety

报告把 helpfulness 和 safety 放在同一个 RL 框架里，而不是把 safety 当成模型发布前的独立补丁。

一个重点是 reward aggregation。很多目标的 reward scale 不同，简单相加会导致大尺度信号压过小尺度信号；而安全这类目标又不能被 response quality 抵消。报告里用了两种策略：

• lexicographic reward shaping：高优先级 reward 打平时，低优先级 reward 才起作用；
• gated reward application：高优先级目标达到最低要求后，低优先级 reward 才应用。

比如安全属于 gated case：unsafe response 直接拿最低 reward，不再因为 response quality 较高而得到补偿。

另一个细节是风格训练。报告里的目标风格包括 warmth without sycophancy、scannable structure、tone calibrated to context、避免冗长开场等。这部分属于 post-training 中面向可用性的训练目标。

Consolidation

三个 specialist 最后要合到一个模型。报告给了 consolidation SFT 的数据 mixture：

这个表显示：general helpfulness / safety 的 sample weight 很高，但 token weight 很低；STEM and Coding 的 token weight 极高。这说明 STEM / code 样本单条更长，尤其 reasoning 和 agentic 轨迹会消耗大量 token。

从 token weight 看，推理模型训练不能只看样本数，还需要看 token budget 和 rollout 长度。

评测

报告里的 STEM 和 agentic coding public benchmark 指标如下。MAI-Thinking-1 的结果是 4 runs average，统一使用 temperature = 1、top-p = 0.97；agentic coding 使用 256K total context length，其他表内 eval 使用 maximum output tokens 256K。

报告中的结果显示，MAI-Thinking-1 在 broad set of benchmark categories 上表现较强，但不是所有任务都领先。比如 Terminal-Bench 2.0 低于 GPT 5.4、Opus 4.6、Kimi K2.6、DeepSeek V4 等；SWE-Bench Pro 接近 Opus 4.6，但低于 GPT 5.4 / Kimi K2.6 / DeepSeek V4 / GLM-5.1。

因此，从报告自身呈现看，MAI-Thinking-1 的定位不是所有 benchmark 的绝对 top-1，而是一个强调自研训练体系、可部署性和持续迭代能力的 reasoning model。

其他 public benchmark 指标如下，Sonnet 4.6 是报告作者在自有 evaluation suite 中生成的结果：

这张表的对比对象比 Table 11 少，原因是很多实验室没有在 model card 或 announcement 里报告这些 benchmark。报告因此只给了 Sonnet 4.6 作为 baseline。

官方介绍页和报告还给出人类 side-by-side evaluation。最终 evaluation set 包含 1276 个英文任务，其中 30% 是 multi-turn。任务来源包括 expert-authored prompts 和经过 PII 过滤的 Microsoft consumer Copilot logs。评估由 Surge AI 管理的 native English raters 完成。

任务分布如下：

Human eval 结果：

整体胜负比例：

也就是说，人类偏好评测中，MAI-Thinking-1 相对 Sonnet 4.6 略占优；相对 Opus 4.6 略落后。分项上，它相对 Sonnet 4.6 的主要优势来自 conciseness/relevance 和 style/tone；instruction following、factuality、completeness 大致在噪声范围内。

安全侧，报告披露了两个内部指标方向：

• safety / over-refusal：低风险 prompt 计算 over-refusal rate，helpfulness 报告为 1 - over-refusal rate；高敏感 item 用 1-5 Likert safety judge，score > 3 视为 safety pass；
• jailbreak：从 vendor、internal red-teaming、HarmBench、StrongREJECT 等来源收集 2.5K unique seed scenarios，再增强成约 9.5K jailbreak prompts；指标是 attack success rate，越低越好。

报告没有在正文文本中给出 safety figure 的逐项数值表，但描述称 MAI-Thinking-1 在 8 个类别中的 5 个相对 Sonnet 4.6 更好或相当，较大提升出现在 CBRN、Self Harm、Elections & Politics；jailbreak ASR 与 Sonnet 4.6 和 Opus 4.6 comparable。

总结与启发

这篇报告的核心关键词是 machine。

报告并不只围绕某个单点技术展开，例如 attention、experts、tokens 或 benchmark 分数；它更强调模型开发体系，即如何把数据、模型、训练、RL、环境和评测组织成持续迭代流程。

这个体系可以拆成几层：

• 数据层：干净、授权、人类生成、可治理；
• 模型层：MoE + local/global attention + LatentMoE，为训练和推理效率服务；
• 实验层：scaling ladder + efficiency gain，避免小规模幻觉；
• 训练层：YOLO、determinism、dropless MoE、goodput；
• RL 层：稳定 GRPO、reward shaping、self-distillation；
• 环境层：SEE、SWE container、tool-use closed-world environment；
• 产品层：helpfulness、安全、style、instruction following；
• 评测层：public benchmark、人类偏好、安全红队。

从训练基础设施视角看，有几个启发比较明确。

第一，训练基础设施需要有明确的目标指标。报告把 goodput 放在生产 KPI 位置，而不只看 MFU 或单步吞吐。对长周期训练来说，节点故障、调度等待、重启、checkpoint 恢复、storage 抖动和通信退化都会进入 wall-clock 成本；因此，基础设施目标应更接近 实际有效训练时间 / 总 wall-clock 时间，而不是单点硬件利用率。

第二，调度系统需要理解硬件拓扑。GB200 这类 rack-scale 系统不是一组同质 GPU 池；NVLink domain、rack boundary、InfiniBand path、reserved spare capacity 都会影响训练稳定性和通信成本。Kubernetes / Kueue / MAI control plane 的价值不只是把 Pod 调起来，而是把 quota、reservation、rack locality 和 topology-aware placement 组织成可持续运行的大作业调度能力。

第三，健康检查和 remediation 要进入调度控制回路。报告中的 certification、NPD conditions、GHR、telemetry、drain、auto remediation 和 recertification 共同决定节点能否进入生产训练池。对训练基础设施来说，坏节点、退化链路和边缘状态存储如果只靠人工排查，会直接降低 goodput；更合理的做法是把健康信号转成可调度状态，并让修复后的节点重新经过 certification。

第四，训练框架和集群控制面需要边界清晰。Kubernetes / Kueue / MAI control plane 负责资源准入、拓扑放置和节点状态；Ray 主要在 admitted job 内部承担 actor 编排和 runtime 管理；YOLO 负责训练循环、sharding、optimizer、checkpoint 和 MoE 相关优化。这个分层可以减少职责混杂，也便于分别优化调度效率、作业恢复和训练性能。

第五，RL 环境越来越接近生产系统。它不是离线数据集，也不是简单 judge，而是一批可执行、可复现、可评分、可并发调度的任务环境。持续构造这样的环境，是获得高质量训练信号的重要条件。因此，训练基础设施不仅要支持 GPU 大作业，也要支持容器化任务环境、工具调用、reward / grader、rollout 和 learner 之间的复杂编排。

第六，可复现和可恢复能力会影响模型迭代速度。报告强调 determinism、checkpoint/restart、dataloader progress、RNG、FP8 scaling history 等状态保存。对于大模型训练，这些能力不只是工程洁癖，而是缩短故障恢复时间、降低实验方差、复盘训练异常和稳定推进 RL climb 的基础条件。

因此，这篇报告可以看成 Microsoft AI 对其自有模型训练体系的一次系统披露：从数据、训练系统、RL 环境到产品评测，形成一个面向 reasoning model 的迭代闭环。

报告仍有不少未披露的部分：比如完整数据来源、各阶段 token / compute budget、更细的 RL rollout 规模、reward model 细节、agentic environment 的实际数量、human eval 的完整 prompt 分布等。这些信息会影响外部读者对 hill-climbing machine 长期效率的判断。

总体看，MAI-Thinking-1 的技术意义不仅在于单次 benchmark 表现，也在于 Microsoft 是否能够把 from-scratch pretraining + 自研 RL infra + 企业场景环境 + Foundry 分发 串成稳定迭代系统。后续 MAI 系列的迭代速度和能力边界，需要继续结合模型发布、评测结果和实际产品表现观察。

MiniMax M2 系列技术报告的标题是 The MiniMax-M2 Series: Mini Activations Unleashing Max Real-World Intelligence。

NVIDIA 也有一篇部署侧性能文章：MiniMax M2.7 Advances Scalable Agentic Workflows on NVIDIA Platforms for Complex AI Applications。

一句话概括：M2 系列不是单纯把模型参数做大，而是在尝试把 低激活 MoE backbone、agent 数据 pipeline、agent-native RL 系统 Forge、长上下文推理优化 组合成一套面向真实 agent workflow 的训练和部署方案。

本文主要整理四类信息：

• 模型参数和结构；
• pre-training / post-training 用了哪些技术；
• Forge 这套 RL 工程系统是什么；
• M2.7 的指标到底提升在哪里。

模型参数

先看 M2 backbone 的基本配置。

这个配置体现出一个明确取舍：总参数很大，但每 token 激活参数控制在 10B 左右。agent task 的 token 消耗通常较高，多轮工具调用、长上下文、观察结果、文件内容都会增加 context 长度。因此，每 token 激活成本会直接影响推理和 RL rollout 成本。

因此，M2 的参数信息需要同时看 total params 和 activated params：229.9B total params / 9.8B activated params。

MoE 结构

M2 的 FFN 层使用 MoE。报告里强调了三个设计：

• fine-grained experts；
• sigmoid gating；
• expert bias。

fine-grained experts 的意思是使用更多、更小的 experts。M2 使用 256 个 experts，每 token 激活 8 个。这样增加了 expert 组合的多样性，也可以降低不同设备之间 expert utilization 的方差。

sigmoid gating 和常见 softmax top-k gating 的区别在于：softmax 有 zero-sum constraint，一个 expert 得分高，其他 expert 得分就会相对被压下去；sigmoid 则是每个 expert 独立打分。报告认为这样可以让多个 expert 同时以较高置信度被激活，routing dynamics 更平滑。

expert bias 则是在 gating score 上加入 per-expert learnable bias，用来改善 load balancing，并减少对 auxiliary load-balancing loss 的依赖。

报告给了一个小规模消融实验：17.8B total params，2B activated params，500B training tokens。

这里 HumanEval 和 MATH 的提升比较明显。

Attention: full attention

M2 使用 full attention，而不是 hybrid SWA、linear attention 或 sparse attention。

这里的 full attention 需要和 causal attention 分开理解。M2 是 decoder-only language model，所以 attention 仍然是 causal 的：当前位置只能看过去和当前位置，不能看未来。full attention 指的是在 causal mask 下，每个 token 可以看完整历史 token，而不是只看一个局部 sliding window。

也就是：

1
2
3
4
5

causal full attention:
token t can attend to token 1 ... t

causal sliding-window attention:
token t can attend to token max(1, t - W) ... t

报告提到，MiniMax 对 hybrid SWA 做了多组继续预训练实验，包括 SWA/full attention 比例、RoPE 设置、layer 内和 layer 间混合、sink token 等。但在 retrieval、多跳推理、in-context learning 和长上下文 agent task 上，SWA variant 有明显损失。

预训练阶段的一组对比：

从这组数据看，32K 内部分任务中 SWA 与 full attention 差距不大；到 128K retrieval 和 long-context ICL 时，局部窗口的覆盖限制开始体现出来。

对于 agent 任务，长上下文通常包含任务描述、工具输出、失败尝试、文件内容、之前的 reasoning state 等信息。如果 attention 机制不能可靠访问完整历史，后续 planning 和 self-correction 会受到影响。

MTP

M2 使用 Multi-Token Prediction。预训练阶段先使用单个 MTP module，K = 1，MTP loss weight 从 0.3 anneal 到 0.1。

在 continued pre-training 的 decay phase，M2 把 MTP module 扩展到 K = 3，用于 multi-step speculative decoding。扩展时不是随机初始化，而是从 main model 复制权重初始化。报告里的解释是：

• copy initialization 收敛更快；
• 随机初始化会带来较高 loss，并短暂干扰 main model；
• 先冻结 main model 训练 MTP modules，loss 稳定后再 joint training。

所以 MTP 在 M2 里有两层作用：

• 训练阶段提供更丰富的预测信号；
• 推理阶段作为 speculative decoding 的 draft path。

这个点和 Forge 也有关。RL rollout 期间 policy 在持续更新，如果 MTP draft model 不跟着适配，acceptance rate 会下降。报告提到 Forge 里 MTP modules 会通过 top-K KL divergence loss 和 RL policy 一起 co-train，从而维持 speculative decoding 的效果。

Pre-training 数据

预训练总量是 29.2T tokens。

报告把它分成：

• constant phase: 19.9T tokens；
• decay phase: 9.3T tokens。

数据来源包括 web documents、academic literature、books、programming code、structured QA。code、math、STEM 会相对自然分布上采样。

长上下文扩展是多阶段完成的：

1

8K -> 32K -> 192K

长上下文数据主要来自：

• high-quality code concatenation；
• naturally long-form PDF documents；
• thematically related document packing。

从报告结构看，M2 的 pre-training 不只服务于通用 base model 能力，也为后续 agent post-training 提供 192K context backbone。

Post-training 数据

报告真正花篇幅的地方其实是 post-training data collection。

M2 系列的 post-training data 不是普通 chat 数据，而是大量带 workspace、tool、environment、verifiable reward 或 artifact-aligned feedback 的 agent trajectories。

主要分几类：

• Agentic Coding；
• Agentic Cowork；
• Reasoning-intensive tasks；
• General conversation and writing；
• Role-play and persona coherence。

但是这里有一个需要注意的点：报告没有披露 post-training 的具体数据规模。

也就是说，论文写了很多 pipeline 和数据类型，但没有给出类似下面这样的硬数字：

• SFT tokens 数；
• RL tokens 数；
• agent trajectories 数；
• SWE/AppDev/Terminal/Cowork 各域样本量；
• rejection sampling 前后 pass rate；
• 每个 stage 的 domain mixing ratio。

报告使用的是 large-scale、corpus、trajectories、at scale 这类描述。能看到的硬规模数字主要集中在 pre-training：29.2T tokens，其中 constant phase 19.9T tokens，decay phase 9.3T tokens。

因此，post-training 部分可确认的信息是：M2 披露了 post-training 数据构造方法和验证信号，但没有披露各类 post-training data 的样本规模和 token budget。

Agentic Coding

Agentic Coding 又分 SWE、AppDev、Terminal-Gym。

SWE pipeline 从 GitHub PR 和 issue 出发，过滤 merged PR、有测试的 PR，再由 agent 构建 Docker environment。之后按 PR 类型做 task routing，比如 bug fix、feature addition、performance optimization、test/refactor 等。

这里最重要的是 reward construction。不同任务类型用不同的验证信号：

• bug fix: F2P / P2P tests；
• feature addition: newly added test points；
• performance optimization: stable and significant performance difference；
• code review: secondary LLM consistency check。

AppDev 则是从零构建应用。它用 expert-in-the-loop 生成 meta queries 和 system prompts，再通过 Agent-as-a-Verifier 做 rejection sampling。AaaV 分三层验证：

• execution layer: 文件、依赖、构建、服务启动、JS error；
• interaction layer: Playwright 检查核心交互；
• visual aesthetics layer: 布局、层次、配色、现代 UI 质量。

Terminal-Gym 则从 Stack Overflow 出发，筛选 terminal-compatible、scriptable、verifiable、Docker-relevant 的任务，然后生成 Dockerfile 和 test script，再做 query evolution 和 difficulty calibration。

这几个 pipeline 的共同点是：数据不只包含最终答案，还包含可执行环境，使结果能够被测试、运行或交互验证。

Agentic Cowork

Cowork 覆盖的东西更接近知识工作者任务：

• deep search and open-web research；
• knowledge-worker office tasks；
• financial analysis and spreadsheet operations；
• slide generation and editing。

这部分的 reward 更复杂。有些任务可以 deterministic check，比如 spreadsheet cell value match、formula recalculation；有些任务需要 rubric-based judge，比如 report、slides、open-ended financial reasoning。

这个方向是 M2.7 相比 M2.5 提升很明显的地方。后面的指标表里，GDPval-AA、MEWC v2、Finance Modeling Pro 都涨得很大。

Reasoning 和 general data

Reasoning data 主要强调 scaling：

• query-side scaling: 扩展题目覆盖；
• response-side scaling: 一个 query 采多个正确解法；
• training-side scaling: 在固定计算预算下调 query expansion / response expansion 比例；
• QA: query、verifier、answer、response 多阶段质量控制。

General conversation and writing 则用于保持通用对话、写作、多轮理解，以及 tool-augmented / tool-free 两种能力。

SFT: interleaved thinking

M2 的 SFT 目标之一是训练 interleaved thinking。

普通 long CoT 往往是：

1

thinking -> final answer

Agent 场景更像：

1

thinking -> tool call -> observation -> thinking -> tool call -> observation -> final answer

报告把它称为 Plan-Act-Reflect loop。关键点是 reasoning state persistence：前一轮的 thinking block 会保留在 history 里，进入下一轮上下文。

这和 stateless per-turn reasoning 不一样。如果每轮工具调用后都把之前的 reasoning 丢掉，模型就要反复重新推导状态，长程任务里容易 drift。

这个地方也能解释为什么 M2 更坚持 full attention：如果 thinking、action、observation 都保留在 192K context 里，模型必须能可靠地访问完整历史。

Forge

是什么

Forge 是 M2 系列的 agent-native RL training system。

它不是模型结构，也不是单个 RL 算法，而是一套让长程 agent trajectories 能进入 RL 训练闭环的工程系统。

普通 RLHF 或 response-level RL 常见输入是：

1

prompt -> response -> reward

Agent RL 的输入则像：

1
2
3
4
5
6
7
8
9

task
-> reasoning
-> tool call
-> observation
-> reasoning
-> tool call
-> observation
-> artifact / test result / final answer
-> reward

轨迹可能有几十 K 到 192K tokens，中间还有工具调用、Docker、browser、spreadsheet、file system、context management、sub-agent delegation。这要求训练系统处理长序列、异步环境和复杂状态转移。

Forge 的系统拆分为：

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

Agent Side
  run agent loop
  manage context
  call tools / envs
  produce trajectories

Middleware
  Gateway Server
  Data Pool

Training / Inference Side
  Rollout Engine
  Train Engine

Agent Side 只负责跑任务和记录轨迹。Rollout Engine 负责高吞吐生成。Train Engine 负责策略更新。Gateway 和 Data Pool 把二者解耦。

一个重要抽象是：Forge 把 LLM generation interface 作为 policy 边界，边界之外的 tool execution、context management、memory access、agent harness 都视为 environment dynamics。

这样它就可以支持两类 agent：

• white-box agent: 训练系统知道它怎么做 context management；
• black-box agent: 训练系统只看到每次发给模型的 state 和模型输出。

black-box support 很重要。因为真实 agent scaffold 往往各有各的 context rewrite、memory、sub-agent、tool protocol。如果要求所有 agent 都白盒接入，系统扩展性会很差。

RL 算法和 reward

M2 系列 RL 使用 CISPO，Clipped Importance Sampling Policy Optimization。报告里给了 objective 和 importance sampling ratio 的 clipped form。

我先不展开公式，抓几个工程上更重要的点：

第一，训练样本是 (state, action) pair。一个 action 是一次 LLM completion，可以包含 reasoning、tool invocation、context operation、sub-agent communication 等。

第二，credit assignment 仍然按完整 episode 做。也就是说，每个 step 的 advantage 要结合整条 trajectory 的结果。

第三，reward 不是只有最终 outcome，而是 composite reward：

1
2
3

r_t = alpha * process_reward_t
    + beta * speed_reward_t
    + performance_reward_t

process reward 用于中间行为，比如工具调用格式、语言混杂、reasoning 结构。

speed reward 用于 wall-clock completion time。这个点很 agent：两个 trajectory 都完成任务，但一个串行慢慢跑，一个并行调用工具，后者对产品更有价值。

performance reward 则是最终任务质量，比如测试通过、artifact 正确、rubric score。

第四，使用 mixed-domain RL。每个 stage 同时混合 reasoning、coding、agent、general 四类数据。这样可以降低单一 agent task RL 造成的 catastrophic forgetting。

Windowed FIFO

Agent rollout 有一个很现实的问题：任务完成时间差异巨大。

1
2
3

simple API task: seconds
coding task: minutes
ML engineering task: hours

如果严格 FIFO，训练会被长任务卡住。
如果谁先完成就训练谁，前期 batch 会被短任务和简单任务主导，后期才出现长任务和难任务，训练分布会漂。

Forge 用 Windowed FIFO 做折中：

1
2

queue = [T0, T1, T2, T3, T4, T5, ...]
window = [T0, T1, T2, T3]

只允许训练系统消费 window 内已经完成的 trajectories。window 内可以乱序取，window 外即使已经完成也不能提前进入训练。

论文里举的窗口大小例子是 W = 0.3N。这个策略牺牲一点点绝对吞吐，换来更稳定的数据分布。

Prefix tree merging

Prefix tree merging 是 Forge 里最有工程味的优化之一。报告称它最高能带来 40x training speedup，并降低显存占用。

Agent RL 训练中，很多 samples 共享长前缀。例如同一个 rollout group 里可能有：

1
2
3

sample 1 = long context + response A
sample 2 = long context + response B
sample 3 = long context + response C

或者同一条 agent trajectory 被拆成多个 step：

1
2
3

s1 -> a1
s2 = s1 + a1 + obs1 -> a2
s3 = s2 + a2 + obs2 -> a3

如果每个 sample 独立 forward，long context 或历史轨迹会被反复计算。

Prefix tree merging 把这些序列组织成一棵 prefix tree：

1
2
3
4

long shared context
├── response A
├── response B
└── response C

或者：

1
2
3
4

context
└── action1 + obs1
    └── action2 + obs2
        └── action3 + obs3

共享 prefix 只 forward 一次。到分叉点后，再分别计算 branch。forward 结束后，根据元数据把 tree 拆回原始 sample，loss 仍然按 sample 独立计算。

它成立的原因是 causal attention：prefix token 的 hidden states 不依赖后续 branch tokens。后面的 token 可以看前面，前面的 token 不会看后面。

概念伪代码：

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

samples = [
    ctx + resp_a,
    ctx + resp_b,
    ctx + resp_c,
]

tree = build_prefix_tree(samples)

def forward_node(node, parent_state):
    state = model_forward_segment(
        tokens=node.tokens,
        parent_state=parent_state,
    )

    for child in node.children:
        forward_node(child, state)

forward_node(tree.root, None)

loss = 0
for sample in samples:
    logits = reconstruct_logits(sample, tree)
    loss += compute_loss(logits, sample.labels)

loss.backward()

真实实现会复杂很多，要处理 attention mask、position ids、loss mask、MoE routing、activation checkpointing、分布式并行和 backward graph。但核心就是把训练 batch 从独立 sequence list 改成共享前缀树。

这个优化和普通 sequence packing 不同。sequence packing 主要减少 padding；prefix tree merging 则是避免重复计算公共历史。对于 192K context 的 agent RL，这类重复计算会带来较高开销。

Rollout 侧推理优化

Forge 还做了几类 inference acceleration。

MTP speculative decoding

M2 的 MTP modules 可以生成 draft tokens，再由 main model 验证。RL 期间 policy 会更新，所以 MTP modules 也要跟着 co-train，否则 draft acceptance rate 会下降。

Prefill-decode disaggregation

把 prefill 和 decode 分开调度。MoE 模型里 prefill 和 decode 的计算形态不同，混在一起容易互相干扰。拆开后可以分别采用更适合的 parallelism 策略。

Global L3 KV cache pool

Agent 多轮交互里有大量共享 prefix。Forge 使用分布式 KV cache pool，提高 prefix cache hit rate。router 会在 queue delay 和 cache migration cost 之间做权衡。

从报告描述看，Forge 里的 rollout engine 不只是离线采样服务，而是包含长上下文、KV cache、MoE serving、speculative decoding、prefill/decode separation、多版本权重同步等能力的推理系统。

训练系统性能

报告明确给出的硬数字是：prefix tree merging 最高可以达到 40x training speedup，同时降低 memory consumption，使更长 sequence 和更大 batch size 成为可能。

其他 Forge 优化更多是定性描述，比如：

• Windowed FIFO 用于在 rollout throughput 和 distributional consistency 之间折中；
• MTP speculative decoding 用于提升 rollout generation throughput；
• prefill-decode disaggregation 用于提升 global throughput 并降低 tail latency；
• global L3 KV cache pool 用于提升 prefix cache hit rate。

但是论文没有给出这些优化各自的 ablation 表，比如没有列出 Windowed FIFO 前后吞吐、GPU utilization、tail latency、KV hit rate、MTP acceptance rate 等数字。

实现架构推断

下面是基于报告描述的工程推断，不是官方源码，也不是论文披露的完整实现。

sequenceDiagram
participant T as Task Queue
participant A as Agent Runner
participant E as Tool / Env Servers
participant G as Gateway Server
participant R as Rollout Engine
participant D as Data Pool
participant Tr as Train Engine

T->>A: task / env spec
loop agent rollout
  A->>G: completion request<br/>state + tools + metadata
  G->>R: normalized request<br/>model version attached
  R-->>G: completion<br/>tokens + logprobs
  G-->>A: action / tool call
  A->>E: execute tool call
  E-->>A: observation / artifact
  A->>D: state / action / observation
  E->>D: verification / reward signal
  R->>D: model version / logprobs
end
D->>D: Windowed FIFO<br/>filtering / batching
D->>Tr: training batch
Tr->>Tr: prefix tree merging<br/>CISPO update
Tr-->>R: updated weights

图里的关键边界是 Gateway：Agent Side 可以保持 scaffold 差异，Training / Inference Side 则通过统一的 completion 接口接收请求、记录元数据，并把 rollout 数据回流到 Data Pool。

Prefix tree merging 可以单独画成下面这个形态：

flowchart LR
b1["before: ctx + a"]
b2["before: ctx + b"]
b3["before: ctx + c"]

ctx["after: shared ctx"]
a["branch a"]
b["branch b"]
c["branch c"]

b1 -. same prefix .-> ctx
b2 -. same prefix .-> ctx
b3 -. same prefix .-> ctx

ctx --> a
ctx --> b
ctx --> c

共享 prefix 只做一次 forward，分叉后的 response segment 分别计算；forward 结束后再根据元数据还原到原始 sample 计算 loss。

根据 Forge 的训练需求，Data Pool 可能需要记录这些字段：

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

trajectory_id
task_id
domain
model_version
states
actions
observations
token_ids
old_logprobs
process_rewards
final_reward
wall_clock_time
tool_calls
artifact_paths
verification_result

Train Engine 的处理流程可以抽象为：

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

trajectories = data_pool.fetch_windowed_fifo_batch()

samples = []
for traj in trajectories:
    for step in traj.steps:
        samples.append({
            "input_ids": step.state_tokens,
            "target_ids": step.action_tokens,
            "old_logprobs": step.old_logprobs,
            "advantage": compute_advantage(traj, step),
        })

batch = prefix_tree_merge(samples)
loss = cispo_loss(batch)

loss.backward()
optimizer.step()

rollout_engine.sync_weights(model)

这里的关键是 old logprobs 和 model version。RL rollout 和 training 之间一定存在 policy lag，所以需要知道 trajectory 是哪个旧 policy 采样出来的，再通过 importance sampling ratio 做修正。

性能数据

M2 系列报告里其余性能数据分布比较散，和 Forge 训练系统性能分开看更清晰。

Agent task 运行设置

这些不是系统吞吐，但可以反映任务成本：

• agent trajectories 最长可到 192K tokens，并可能包含 thousands of intermediate actions；
• rollout completion time 从 seconds 到 hours；
• Terminal-Bench 2.0 使用 8 vCPU / 16GB sandbox，2 小时 wall-clock timeout，4 trials；
• MLE Bench Lite 对 22 个 competitions 运行，每个 competition 在 single-A30 sandbox 中跑 24 小时，最终取 3 个 independent 24-hour trials 的平均 medal rate；
• VIBE-Pro、HyperTask、MM Claw、MEWC v2、Finance Modeling Pro 等多项 agent / artifact benchmark 使用 3 trials。

Self-evolution 内部效率

• M2.7 在 RL team workflow 中吸收 30% 到 50% 的 daily iteration workload；
• 对内部 programming scaffold 做 100-round autonomous iteration；
• 引入 loop detection 和更好的参数组合后，内部评估有 30% performance gain。

这部分属于内部系统和内部评测，不是外部可复现 benchmark。

部署侧推理性能

来自 NVIDIA 技术博客，而不是 M2 论文主体。NVIDIA 提到在 Blackwell Ultra GPU 上，针对 MiniMax M2 系列在 vLLM / SGLang 集成 QK RMSNorm kernel 和 FP8 MoE kernel 后，在 1K/1K ISL/OSL dataset 上：

• vLLM throughput 最高提升 2.5x；
• SGLang throughput 最高提升 2.7x。

这个数据属于部署工程部分，不应和 Forge 训练系统性能混在一起。Forge 是 post-training / RL infrastructure；NVIDIA 这里讲的是 open-source inference framework 的 serving optimization。

指标

论文 Table 4 给了 M2.7、M2.5 和几个闭源 frontier baseline 的对比。这里只摘 M2.7 和 M2.5。

从表中可以看到：

• M2.7 相比 M2.5 的大幅提升集中在 agent / cowork / office / MLE；
• Finance Modeling Pro 从 33.8 到 57.0；
• GDPval-AA 从 35.0 到 50.0；
• MEWC v2 从 49.8 到 63.3；
• MLE Bench Lite 从 51.5 到 66.6；
• MMLU-Pro 从 85.2 降到 81.8。

这说明 M2.7 不是所有传统静态知识 benchmark 都提升。报告更强调的是：agent data pipeline 和 Forge RL 对真实 workflow benchmark 的提升。

M2.7 的 self-evolution

报告里还提到 M2.7 的 self-evolution。

MiniMax 的说法是，M2.7 可以在内部 Model Iteration System 里帮助 RL 团队：

• profile ongoing runs；
• read logs；
• diagnose metric anomalies；
• debug code；
• adjust configs；
• generate reports；
• modify agent scaffold。

报告称它可以吸收 RL 团队日常 30% 到 50% 的 iteration workload。另一个例子是，M2.7 对内部 programming scaffold 做了 100-round autonomous iteration，引入 loop detection 和更好的参数组合，在内部评估上带来 30% performance gain。

这部分高度依赖内部工作流和内部评测，应视为官方披露的内部案例，而不是外部可复现实验结论。

总结

M2 系列报告的重点不只是单个 benchmark，而是一套完整路线：

1
2
3
4
5
6
7

low-activation MoE backbone
-> long-context full attention
-> MTP for training signal and speculative decoding
-> verifiable agent trajectory data
-> interleaved thinking SFT
-> Forge agent-native RL
-> rollout / training / serving co-optimization

如果只看参数，M2 是一个 229.9B total / 9.8B activated 的 MoE 模型。
如果看训练，它是一个围绕可验证 agent trajectories 做 post-training 的模型。
如果看工程，Forge 才是这篇报告里很关键的东西：它把 agent loop、推理服务、轨迹存储、reward、RL trainer 和权重同步接成一个系统。

这也是 M2 系列和很多只讲模型结构的技术报告不同的地方。它把模型能力放在完整 agent workflow 里讲，重点不是“模型会不会回答”，而是“模型能不能在环境里把事情做完，并且这个训练闭环能不能规模化”。

Ray 默认把 object 放进 Plasma object store——每个节点上一个基于 共享内存 的本地 store 进程，worker 通过它读写 object。ray.put()、task / actor 返回值等先落本机 Plasma；跨节点时再由 Ray 的 ownership / scheduling 层协调 fetch，但 各节点本地内存层仍是 Plasma。内存不够时会 spill 到磁盘（默认在 session 临时目录下），需要时再 restore 回 Plasma；ray memory 里的 Plasma memory usage 就是这一层。

task / actor 消费 object 时要反序列化。对 CUDA torch.Tensor 来说，默认路径意味着 GPU → CPU（进 Plasma）→ GPU 的来回拷贝，在 actor 间频繁传 tensor 时开销很大。

graph LR
subgraph N1["节点 1"]
  A1["Producer Actor<br/>GPU tensor"]
  B1["Plasma Store"]
end
subgraph N2["节点 2"]
  A2["Consumer Actor"]
  B2["Plasma Store"]
end
A1 -->|"GPU to CPU, 入 Plasma"| B1
B1 -->|"跨节点 fetch"| B2
B2 -->|"拷贝到 GPU"| A2
A1 -->|"RDT 使用 Gloo/NCCL/NIXL<br/>在 actor 间 send/recv"| A2
style B1 fill:#eee,stroke:#ccc,color:#999
style B2 fill:#eee,stroke:#ccc,color:#999
style A1 fill:#ddeeff,stroke:#338
style A2 fill:#ddeeff,stroke:#338

上：传统 Plasma 路径需多次 CPU/GPU/内存拷贝；下：RDT 经 Gloo/NCCL/NIXL 在 actor 间 send/recv，绕开 Plasma。

Ray Direct Transport (RDT) 是在 ObjectRef 语义上做的增强：tensor 留在 producer actor 侧（GPU 上），consumer 需要时由 Ray 协调两端做 send/recv，绕开 Plasma object store 的序列化与拷贝。底层可选 Gloo / NCCL / NIXL——Gloo、NCCL 是 collective 库，需先建 collective group，再在 group 内走 p2p 传输；NIXL 则是基于 UCX 的 p2p RDMA，无需预建 group，且 ray.get 可走 one-sided 取回。

RDT 目前仍是 alpha，API 和限制都可能变；下文基于 Ray 2.55 文档整理。

基本用法

在返回 torch.Tensor 的 actor method 上加 @ray.method(tensor_transport=...)：

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

import torch
import ray
from ray.experimental.collective import create_collective_group

@ray.remote
class MyActor:
    @ray.method(tensor_transport="gloo")
    def random_tensor(self):
        return torch.randn(1000, 1000)

    def sum(self, tensor: torch.Tensor):
        return torch.sum(tensor)

sender, receiver = MyActor.remote(), MyActor.remote()
group = create_collective_group([sender, receiver], backend="torch_gloo")

tensor = sender.random_tensor.remote()
result = receiver.sum.remote(tensor)
print(ray.get(result))

• decorator 只加在产出 tensor 的方法上，消费方不用加（除非它也要返回 RDT tensor）。
• tensor 存在 producer actor 里，不是 Plasma object store。
• 传给另一个 actor 时，Ray 自动用指定 transport 做 send/recv。
• 返回值若未标注 RDT，仍走默认 Plasma object store（上例 sum 的标量结果）。

嵌套结构、多 tensor 返回值也支持，Ray 会递归识别其中的 torch.Tensor。

三种 transport

Gloo / NCCL 是 collective 语义，使用前必须 create_collective_group，且 backend 与 tensor_transport 一致。NIXL 更灵活，actor 环境装好 nixl 即可，适合跨节点 p2p。

NCCL 版几乎就是 Gloo 版三处替换：tensor_transport="nccl"、backend="nccl"、tensor 放 GPU。

NIXL 额外支持 driver 侧 ray.put(t, _tensor_transport="nixl")，以及 consumer 内 ray.get(ref) 直接经 NIXL 取回。

collective transport 的 ray.get 若 caller 不在 group 里会报错，需配置 _use_object_store=True 回退。

与 Plasma object store 的语义差异

RDT object 是可变的。 Ray 只持有 tensor 引用，不做 immutable copy。producer 若仍持有同一块 tensor 并在 in-place 修改，后续 consumer 可能看到被改过的数据。这与 Ray Core 默认「actor 返回即拷贝」的行为不同。

传回 同一个 producer actor 时零拷贝，只是引用；若同时再传给别的 actor，in-place 修改会影响 Ray 内部持有的那份，Ray 会打印 warning。

需要 producer 再次写同一块 tensor 时，用 ray.experimental.wait_tensor_freed(tensor) 等 Ray 释放所有引用；注意此时 driver 不要再 ray.get 持有该 ref，否则会死锁。

限制

当前 alpha 状态

• 仅 torch.Tensor，仅 Ray actor（不含普通 task）。
• 不支持 asyncio（tracking issue）。
• Gloo / NCCL：
  • 只有 创建 collective group 的进程 能提交返回 / 传递 RDT object 的 actor task。
  • RDT ObjectRef 不能序列化后跨进程传递，只能作为 同 group 内 actor task 的直接参数。
  • 每个 actor 在同一 transport 下同时只能属于一个 group。
  • 不支持 ray.put。
• NIXL：同一 actor 上若先后存两个 object、tensor 集合有重叠但不完全相同，当前有已知问题；需等第一个 ObjectRef 出 scope 后再存第二个。

系统级传输错误：Gloo/NCCL collective 失败会 销毁 group 并 kill actor；NIXL 会 abort 并在依赖 task / ray.get 处抛异常。超时可调 RAY_rdt_fetch_fail_timeout_milliseconds。

与 RL 训推 infra 的关系

RL 里 actor 间传 rollout buffer、logits、hidden states 若走默认 Plasma object store，GPU 数据会被反复拉到 CPU。RDT 把这条路径收成 actor 间 direct transport，和 NCCL collective、NIXL RDMA 对齐，适合 多 actor 流水线（例如 rollout actor → trainer actor）且 tensor 较大的场景。但 alpha 阶段的 collective group 创建进程限制、可变语义、以及仅 actor 支持，使用前要先对照 workload 评估是否适用。

参考

• Ray Object Spilling（Plasma 内存层与 spill）
• Ray Direct Transport 官方文档
• NIXL

想着在这个时间点上回顾下模型和模型训练 infra 发展的经历，就以 DS 的技术报告为例吧。

DS 系列模型：

1. DeepSeek LLM — DeepSeek LLM 7B、DeepSeek LLM 67B Dense（2024/01/05） arXiv:2401.02954
2. DeepSeek-Coder — 1.3B / 6.7B / 33B（2024/01/25） arXiv:2401.14196
3. DeepSeekMoE — MoE 语言模型系列 2B / 16B / 145B 等（2024/01/11） arXiv:2401.06066
4. DeepSeekMath — DeepSeekMath-7B（2024/02/05） arXiv:2402.03300
5. DeepSeek-V2 — 第二代 MoE 通用大模型（2024/05/07） arXiv:2405.04434
6. DeepSeek-V2.5 — 通用与代码能力合流迭代（2024/09/06；无单独 arXiv 技术报告，架构见 V2） 官方说明
7. DeepSeek-V3 — 第三代 MoE 通用大模型（2024/12/27） arXiv:2412.19437
8. DeepSeek-R1、DeepSeek-R1-Zero（2025/01/22） arXiv:2501.12948

继而再看框架的实现。不是之前不能做，之前这是个浩大的工程。现在借助模型不仅是代码门槛下降了，理解 sota 工作的门槛也下降了，可以抽空广泛的了解起来了，成为新时代的 “全栈” 工程师。

DeepSeek LLM

LLaMA

2T tokens pre train

1M sft, RLHF, SFT → RM → PPO 这条经典 RLHF pipeline

DPO, 不显式训一个单独的 RM、也不做 RL 循环，直接用 偏好对 数据（同一条 prompt 下，人类更喜欢回答 A 而不是 B）去更新语言模型。

SFT -> DPO

模型架构

a Pre-Norm structure with RMSNorm (Zhang and Sennrich, 2019) function

using SwiGLU (Shazeer, 2020) as the activation function

Rotary Embedding (Su et al., 2024) for positional encoding

Grouped Query Attention (GQA)

AdamW optimizer (Loshchilov and Hutter, 2017), with the following hyperparameters: 𝛽1 = 0.9, 𝛽2 = 0.95, and weight_decay = 0.1

HAI-LLM = Megatron 式多并行 + Flash Attention + ZeRO-1 省优化器显存 + 通信计算重叠 + 融合 kernel；数值上用 bf16 算、fp32 攒梯度保稳；最后在 softmax/CE 上用 in-place 省 logits 显存。

DeepSeek-Coder

代码生成和代码补全，专门进行 FIM 代码补全训练

employ HuggingFace Tokenizer library 使用 Byte Pair Encoding 技术在训练语料子集上进行训练得到

基于 DeepSeek LLM 模型架构和训练技术，decode-only transformer, RoPE, GQA, FlashAttention v2

AdamW 优化器

并行策略实现仍然用的自研的 HAI-LLM 框架

DeepSeek-Coder 主系列（1.3B / 6.7B / 33B）：from scratch pt 2T -> Base（技术报告：pre-training 含 FIM、context 扩至 16K；DeepSeek-Coder 仓库 README Model Training 两阶段 1.8T@4K + 200B@16K）；SFT 2B -> Instruct

DeepSeek-Coder-v1.5（仅 7B）：CPT from DeepSeek-LLM 7B，2T@4K（next-token，无 FIM/16K）

DeepSeekMoE

更小（多）的专家 + 共享专家

2B 实验

DeepSeekMath

DeepSeek-V2

DeepSeek-V2.5

DeepSeek-V3

DeepSeek-R1

1. PPO

Proximal Policy Optimization Algorithms

两阶段循环

为什么可以“多轮”

通常情况下，如果对同一批数据进行多轮优化，策略会因为更新过头而崩溃。但 PPO 引入了 Clipped Objective（裁剪目标函数）：

• 安全护栏：在每一轮优化中，PPO 会计算新策略和采样时的旧策略的概率比。如果这个比值超出了设定的范围（比如 $0.8 \sim 1.2$），梯度就会被“截断”。
• 效果：这确保了即使在这一批数据上反复“薅羊毛”优化，新策略也不会跑得离旧策略太远，从而保证了训练的稳定性。

1.1. 采样阶段 (Sampling Phase)

• 动作：让当前的策略 $\pi_{\theta_{old}}$ 在环境中运行一段时间。
• 产出：收集一批轨迹数据（包括状态 $s$、动作 $a$、奖励 $r$ 等）。
• 性质：这些数据是“新鲜”的，反映了当前策略的行为模式。

在这个阶段，神经网络的参数是固定不动的（即 $\theta_{old}$）。Actor (策略网络)：在环境中根据概率分布选择动作。数据收集：把 $(s_t, a_t, r_t, s_{t+1})$ 存入一个临时的 Buffer。目标：收集足够数量的轨迹（比如 2048 个时间步）。

1.1.1. 计算“标签” (Preprocessing)

在开始训练前，利用收集到的数据计算两个关键值：

• $\hat{A}_t$ (Advantage)：优势函数，用来衡量这个动作比平均水平好多少。
• $R_t$ (Returns)：这一步动作带来的累积奖励。

注意到

• $r_t(\theta)$：新旧策略概率比（用于 Actor）。
• $\hat{A}_t$：优势估计（用于 Actor，决定更新方向）。
• $R_t$：回报目标值（用于 Critic，提升估值精度）。

如果只用即时奖励 $r_t$ 作为目标，Critic 就会变得非常“短视”。即时奖励 $r_t$：只代表当前这一步的好坏。回报目标 $R_t$：代表从当前时刻起，在轨迹剩余部分上累计（折现）后的总回报。目标是让 Critic 具备“向前看”的能力，因此用 $V(s_t)$ 去拟合这个 $R_t$。

$$R_t = \hat{A}_t + V(s_t)$$

$R_t$ (Returns)：作为 Critic 网络的监督信号（标签）。

• 计算逻辑：通过 $\hat{A}_t$（优势）与采样时旧的 $V(s_t)$ 相加得到：$R_t = \hat{A}_t + V(s_t)$。
• 物理意义：它代表了在当前策略下，从状态 $s_t$ 开始预期能获得的折现总奖励。Critic 的优化目标就是让预测值 $V_\theta(s_t)$ 尽可能接近这个 $R_t$。

这意味着：

1. 先用 GAE 算出了优势估计 $\hat{A}_t$。
2. 通过 $\hat{A}_t + V(s_t)$，便可反向推导出这一步动作对应的“目标回报” $R_t$。
3. 价值损失 (Value Loss) 就变成了：$MSE(V_{new}(s_t), R_t)$。

1.1.2. 小结

• Actor：利用 $\hat{A}_t$（相对好坏）来决定 $\theta$ 的更新方向。
• Critic：利用 $R_t$（绝对得分）来修正自己对世界的认知。

1.2. 优化阶段 (Optimization Phase)

多轮优化 (Several Epochs)：

• 动作：将刚才采样的这一批数据反复输入神经网络进行多次梯度更新。
• 关键点：在传统的 On-policy 算法（如普通的策略梯度）中，这批数据更新一次就必须扔掉。但 PPO 允许在同一批数据上跑 3 轮、5 轮甚至 10 轮（Epochs）。

要把 Buffer 里的数据，分成更小的 Mini-batches，重复训练 $K$ 个 Epochs（比如 $K=10$）。在每一个 Epoch 里的微观操作：计算概率比 $r_t(\theta)$：用当前正在更新的 $\theta$ 计算动作概率，除以采样时的 $\theta_{old}$ 计算的概率。应用裁剪 $CLIP$：如果 $r_t(\theta)$ 偏离 1 太远（比如超过 20%），就强行截断。梯度更新：通过反向传播更新参数 $\theta$。

为什么 $r_t(\theta)$ 允许“多轮更新”。PPO 能够从 On-policy 转向近乎 Off-policy 的理论支柱，PPO 本质上是利用了重要性采样技术。

• 理论背景：优化目标是新策略 $\pi_\theta$，但训练数据来自旧策略 $\pi_{\theta_{old}}$ 的采样分布。
• 补偿机制：通过概率比率 $r_t(\theta)$，对数据分布偏差做重要性采样修正。
• 约束：重要性采样要求两个分布不能差太远，否则方差会爆炸。这正是 $L^{CLIP}$ 存在的根本原因——它在数学上维护了重要性采样的有效区间。

1.2.1. 优势估计

通常采用 GAE (Generalized Advantage Estimation)。

简单来说，优势函数 $\hat{A}_t$ 的目标是回答：“在状态 $s_t$ 下采取动作 $a_t$，比平均情况（即 Baseline）好多少？”