本文是《通信与互联：从 NCCL 到 RDMA》系列的第 6 篇（共九篇）。上一篇：PyTorch 的通信栈——ProcessGroupNCCL、stream 语义与计算通信重叠；下一篇：推理侧的通信——custom all-reduce 与 KV 传输。

前五篇建立了一条完整的因果链：第一篇的 α-β 模型给出一次集合通信的理论时间，第二、三篇给出链路能提供的 β 和 α，第四篇讲 NCCL 如何在探测到的拓扑上选 ring/tree、channel 数、算法与协议去逼近这个上限，第五篇讲 ProcessGroupNCCL 如何把 NCCL kernel 放到自己的 stream 上、watchdog 如何盯着每一个 WorkNCCL。这条链上每一环都可能出错，而错误的表现只有三种：慢、卡、结果不对。本篇的任务是把前五篇变成一套可操作的方法——面对这三种现象，先测什么、先看什么、先改什么。

性能这一半的核心工具是 nccl-tests。它做的事情很简单：对一组消息大小依次调用同一个集合通信原语，报告时间、algbw 和 busbw。但这条曲线包含的信息远比”峰值多少 GB/s”多：小消息端的平台高度是 α，大消息端的平台高度是 β，拐点位置是二者之比，中段的每一次台阶都是算法或协议的切换点。学会读这条曲线，再配合 NCCL_DEBUG=INFO 里的几行决策日志，绝大多数”通信慢”的问题都能在一小时内定位到某一层。

正确性这一半的核心问题是 hang。集合通信有一个和普通 RPC 完全不同的失败模式：它没有”对端拒绝”或”连接断开”这样的即时错误，一个 rank 少调用一次 all_reduce、或者用了不同大小的 tensor，其他所有 rank 只会在 GPU 上安静地自旋等待，直到 watchdog 的 timeout 到期。到那时每个 rank 的日志都停在同一个地方，看起来所有人都有责任，而实际上只有一个 rank 掉队了。找到它需要把所有 rank 最近的集合通信记录按序号对齐——这正是 PyTorch Flight Recorder 做的事，NCCL 2.24 之后自带的 RAS 子系统提供了另一种视角。

本篇要回答总纲提出的核心问题：

一个 64 卡训练任务在第 3000 步 hang 住，所有 rank 的日志都停在 all_reduce。是谁的问题、是哪一次 all_reduce、为什么会等到 timeout 才暴露？¹

依照系列惯例，本文的性能数字要么是可推导的理论值，要么是公开标称值，要么是”通常能达到”的区间，全部明确标注为非实测；曲线形状是典型形状，不是某台机器的测量结果。源码以 NCCL 2.28.9、nccl-tests 2.18.3 与 PyTorch 2.12（v2.12.0）为准；TORCH_NCCL_* 环境变量的默认值随版本变过，正文只标注 PyTorch 2.12 源码里的默认值，读者以自己手上的版本为准。

一、总览：从现象到层级的排障框架

1. 两半、三种现象、一张层级图

本篇分两半：性能这一半（第二到五章）回答”慢”——用 nccl-tests 测、用日志看决策、用参数调；正确性这一半（第六到八章）回答”卡”与”错”——hang 的分类与定位、timeout 的语义、数值问题。两半面对的是同样的六层，所以先把层级图放在这里。

一次 dist.all_reduce 从 Python 到线路要穿过六层，每一层都有自己的观测手段和典型故障：

三种现象在这张图上的排查方向不同。慢从下往上：先确认硬件和网络能提供的上限（第二、三篇的 nvbandwidth 与 ib_write_bw），再用 nccl-tests 看 NCCL 拿到了多少，最后看框架侧是否浪费了。卡从上往下：先用 Flight Recorder 判断是不是调用不一致，是则到此为止；不是再往下看 rank 是否存活、网络是否断开。错几乎只在框架层和数值层：归约顺序、算法差异、stream 竞争、NaN 传播。

2. 四段法在本篇的落法

总纲规定的”算一算、看一看、测一测、比一比”在本篇是主线而不是点缀：

• 算一算：用第一篇的 ring 公式预测 busbw 曲线的形状与拐点，算出”这条曲线应该长什么样”；
• 看一看：读 NCCL_DEBUG=INFO 日志里 transport、算法、协议、channel 的决策，读 src/misc/param.cc 和 src/debug.cc 弄清环境变量与日志的机制；
• 测一测：nccl-tests 扫描，nccl-tests 之外还有 py-spy、gdb、Flight Recorder 这些”测卡在哪”的工具；
• 比一比：实测曲线与预测曲线的每一处偏差都对应一个检查项，本篇第九章把它们整理成决策树。

3. 本文的章节安排

二、nccl-tests：怎么测

1. 编译与二进制

nccl-tests 是 NCCL 官方的性能与正确性测试集，每个集合通信原语一个二进制：all_reduce_perf、all_gather_perf、reduce_scatter_perf、broadcast_perf、reduce_perf、alltoall_perf、sendrecv_perf、scatter_perf、gather_perf、hypercube_perf，对应 src/*.cu 里的同名文件；公共逻辑（参数解析、buffer 分配、计时、输出）在 src/common.cu 与 src/util.cu，数据校验的生成与比对在 verifiable/verifiable.cu。二进制启动时会打一行 # nccl-tests version 2.18.3 nccl-headers=<编译时 NCCL 版本码> nccl-library=<运行时 NCCL 版本码>（src/util.cu，版本号来自 src/common.h 的 NCCL_TESTS_VERSION）。

编译分两种。不带 MPI 只能单进程：

make -j CUDA_HOME=/usr/local/cuda NCCL_HOME=/path/to/nccl

带 MPI 才能多进程、多机（README 明确说明多进程数由 MPI 管理，不作为参数传入）：

make -j MPI=1 MPI_HOME=/path/to/openmpi CUDA_HOME=/usr/local/cuda NCCL_HOME=/path/to/nccl

NCCL_HOME 决定链接哪个 libnccl.so。运行时 LD_LIBRARY_PATH 若指向另一个版本，测出来的是另一个版本的数字——排障时第一件事是 ldd ./build/all_reduce_perf | grep nccl 确认实际加载的库，并核对 nccl-tests 开头那行 nccl-headers=… nccl-library=… 与 NCCL 日志里的 NCCL version 2.x.y。

2. 参数

src/common.cu 的 getopt_long 表定义了全部参数（默认值见同文件开头的全局变量与 README 的 Arguments 一节），常用的这些：

-b, --minbytes        起始消息大小（可带 K/M/G 后缀，按 1024 进位；默认 32M）
-e, --maxbytes        结束消息大小（默认 32M）
-f, --stepfactor      每步乘以的因子（-f 2 即翻倍扫描）
-i, --stepbytes       每步加的字节数（与 -f 二选一；默认 1M）
-g, --ngpus           每个线程（进程）管几张 GPU
-t, --nthreads        每进程几个线程（每线程再管 -g 张卡；单机不用 MPI 时通常 -t 1 -g 8）
-n, --iters           每个消息大小计时迭代次数（默认 20）
-w, --warmup_iters    预热迭代次数（默认 1，几乎总要手动加大）
-m, --agg_iters       每次迭代聚合几个操作（ncclGroupStart/End 包起来；默认 1）
-N, --run_cycles      整个扫描重复几轮并逐轮打印（默认 1；0 无限）
-c, --check           校验迭代次数；-c 0 关闭校验（默认 1，即校验一次）
-o, --op              归约算子：sum/prod/min/max/avg/mulsum/all
-d, --datatype        数据类型：int8/uint8/int32/uint32/int64/uint64/half/float/double/bfloat16/…/all
-r, --root            broadcast/reduce 的 root
-z, --blocking        0 不阻塞（默认，计时后再同步）；1 每次调用后 wait + barrier；2 每次调用后 wait 但不 barrier
-a, --average         多进程时报告哪个时间：0=rank0 / 1=平均（默认）/ 2=最小 / 3=最大
-G, --cudagraph       把 -n 次迭代捕获进 CUDA Graph 再 launch <num> 次，去掉 CPU 侧 launch 开销
-R, --local_register  1 本地注册 send/recv buffer；2 对称注册（NCCL 2.27+ 的 symmetric memory 路径）
-x, --cta_policy      0/1/2 设置 NCCL_CTA_POLICY（NCCL 2.27+；2 需 2.28+）
-C, --report_cputime  1 时报告 CPU 侧时间而不是 GPU 侧延迟
-S, --report_timestamps  每行结果加时间戳
-J, --output_file     结果写 JSON
-T, --timeout         秒；超过则报错退出，而不是永远 hang
-p, --parallel_init   1 时用多线程并行初始化 communicator
-M, --memory_report   1 时结尾打印显存占用报告
-u, --unalign         把 send/recv buffer 故意偏移若干元素，测非对齐路径

总 rank 数等于 进程数 × -t × -g。-g 8 单进程管 8 卡是最简单的节点内测试；多机则每进程一卡 -g 1，让 MPI 起 8 × 节点数个进程。nccl-tests 2.18.3 的结果表里没有算法/协议/channel 数这几列；要知道第四篇讲的”NCCL 为这个消息大小选了什么”，用 NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=TUNING 让 rank 0 在每次调用时打一行 AllReduce: N Bytes -> Algo Ring proto LL128 channel{Lo..Hi}={0..15}（NCCL src/enqueue.cc），再与结果表按消息大小对照。

3. 三种启动方式

单机不用 MPI：

./build/all_reduce_perf -b 8 -e 8G -f 2 -g 8 -n 50 -w 10 -c 0

多机用 mpirun（-N 8 每节点 8 进程，-np 总进程数；-x 把环境变量传过去）：

mpirun -np 16 -N 8 --hostfile hosts \
  -x NCCL_DEBUG=INFO -x NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_4,mlx5_5,mlx5_6,mlx5_7 \
  -x NCCL_SOCKET_IFNAME=eth0 -x LD_LIBRARY_PATH \
  ./build/all_reduce_perf -b 8 -e 8G -f 2 -g 1 -n 50 -w 10 -c 0

src/common.cu 的 run() 里可以看到 MPI 只用来做几件事：MPI_Comm_rank/size 确定 rank、MPI_Comm_split 按 NCCL_TESTS_SPLIT 的 color 分组、MPI_Bcast 分发 ncclUniqueId、结果汇总（Allreduce() 里的 MPI_Allreduce 取时间的平均/最小/最大，以及错误数求和）。数据路径与 MPI 无关，所以 Slurm 下 srun --mpi=pmix 也可以。localRank 由 hostname 的哈希（MPI_Allgather 收集 hostHashs）算出，用它选 GPU；NCCL_TESTS_DEVICE 可以覆盖起始 GPU 编号。

nccl-tests 本身没有 torchrun 模式。但很多集群没有 MPI，只有 PyTorch，这时等价的做法是一个几十行的 PyTorch 脚本在 torchrun 下做同样的扫描——第十章的 sweep.sh 附带了这个脚本（torch_sweep.py）。要注意它测的是”经过 ProcessGroupNCCL 的 NCCL”，多了第五篇讲的 event 记录、watchdog 与 stream 切换的开销，小消息端会比 nccl-tests 高几微秒，大消息端应当一致；如果大消息端也差很多，说明框架侧有问题（比如 TORCH_NCCL_HIGH_PRIORITY、stream 优先级、或者 tensor 不连续触发了拷贝）。

4. 输出的每一列

一次扫描的输出长这样（src/util.cu 的 writeResultHeader 决定列名，writeResultFooter 打结尾三行）：

# nccl-tests version 2.18.3 nccl-headers=22809 nccl-library=22809
# Using devices
#  Rank  0 Group  0 Pid  12345 on   hostname device  0 [0000:1a:00.0] NVIDIA H100 80GB HBM3
   ...
#
#                                                              out-of-place                       in-place
#       size         count      type   redop    root     time   algbw   busbw #wrong     time   algbw   busbw #wrong
#        (B)    (elements)                               (us)  (GB/s)  (GB/s)            (us)  (GB/s)  (GB/s)
           8             2     float     sum      -1    xx.xx    0.00    0.00      0    xx.xx    0.00    0.00      0
   ...
  1073741824     268435456     float     sum      -1  xxxx.x   xxx.x   xxx.x      0  xxxx.x   xxx.x   xxx.x      0
# Out of bounds values : 0 OK
# Avg bus bandwidth    : xxx.xx
#
# Collective test concluded: all_reduce_perf

• size 是字节，count 是元素数；
• time 是 -n 次迭代的平均单次时间（微秒），多进程时按 -a 取 rank 0 / 平均 / 最小 / 最大；
• algbw = size / time，busbw = algbw × 校正系数，系数在每个 src/*.cu 的 *GetBw 里：AllReduceGetBw 乘 \(\frac{2(n-1)}{n}\)，AllGatherGetBw、ReduceScatterGetBw、AlltoAllGetBw 乘 \(\frac{n-1}{n}\)，BroadcastGetBw、SendRecvGetBw 乘 1；
• #wrong 是校验发现的错误元素数，-c 0 时为 0 且不做校验；
• out-of-place 与 in-place 是 send/recv buffer 不同与相同两种情况，正常两组数字接近；
• 最后的 Avg bus bandwidth 是所有大小、out-of-place 与 in-place 两组 busbw 的算术平均（common.cu 里每打一行结果就 bw[0] += busBw; bw_count[0]++），只是一个粗略的回归比较值，不能当峰值用；NCCL_TESTS_MIN_BW 环境变量设一个 GB/s 阈值，平均值低于阈值的 90% 时该行标 FAILED 并返回非零退出码，用于 CI 验收。

busbw 是唯一能与硬件带宽直接比较的数字，理由第一篇讲过：ring all_reduce 每个 rank 实际收发 \(\frac{2(n-1)}{n}S\) 字节，busbw 就是这个流量除以时间，也就是”链路实际被驱动到了多少”。

5. 测量的纪律

几条不遵守就得不到可比数字的规则：

• 预热要够：-w 至少 5～10 次（默认只有 1 次）。第一次调用触发 kernel 加载、connection 建立、buffer 注册，可能是稳态的几十倍；
• 迭代要够、又不能过长：小消息 -n 100，大消息 -n 20 已足够；8 GB 消息 100 次迭代要跑几分钟；
• 校验单独跑：-c 1（默认）会在每个大小上、计时迭代之外额外跑一次带校验的迭代，并用 ncclVerifiableVerify kernel 在 GPU 上与预期值比对，不影响计时数字但拖慢整体；性能扫描用 -c 0，正确性检查再用 -c 1 -o all -d all；
• 控制干扰：其他 GPU 任务、其他网络流量、CPU 频率策略都会影响小消息端；多机测试前先各节点跑一遍单机基线；
• 固定并记录环境：NCCL 版本、NCCL_* 环境变量、驱动、固件、nvidia-smi topo -m 的输出与结果一起存档，否则三个月后无法解释两条曲线为什么不同；
• 加 -T：多机测试加 -T 600，网络有问题时得到一个错误而不是一个永远不退出的进程。

三、带宽曲线的读法

1. 横轴、纵轴、两个平台与拐点

把 size 取对数作横轴、busbw 作纵轴，一条健康的 all_reduce 曲线是一个 S 形：

busbw
  ▲
  │                                            ┌──────────────  带宽平台：≈ β_eff
  │                                        ┌───┘
  │                                    ┌───┘
  │                                ┌───┘  ← 拐点区：α 项与 β 项相当
  │                            ┌───┘
  │                        ┌───┘
  │                    ┌───┘
  │              ┌─────┘
  │  ────────────┘  ← 延迟平台：busbw ∝ S，斜率 45°（对数坐标下），高度由 α 决定
  └────────────────────────────────────────────────────────────▶ size (log)
     8B    1K    32K   1M    4M   16M   64M  256M  1G   8G

左端 busbw 随 S 线性增长——因为时间几乎不变（被 α 主导），busbw = 常数系数 × S / T。在对数坐标下这是一条 45° 斜线，它的纵向位置由 α 决定：同样 1 MB 的消息，α 从 10 µs 变到 20 µs，这段曲线整体下移一半。看小消息端时应该直接看 time 列而不是 busbw 列：8 字节 all_reduce 的 time 就是这套系统的 α 下限（kernel 启动 + 一轮握手）。

右端 busbw 趋于平台——时间被 S/β 主导，busbw 趋于链路的有效带宽 β_eff。平台高度与硬件标称值的比值就是”效率”，通常能达到的区间在 80～95%。

中间是拐点区。第一篇定义了拐点：α 项与 β 项相等的消息大小。对 ring all_reduce：

在 \(S_{\text{knee}}\) 处 busbw 恰好是平台的一半；要到平台的 90%，带宽项须是延迟项的 9 倍，即 \(S \approx 9\,S_{\text{knee}}\)。

2. 用 α-β 模型预测拐点

代入典型值（非实测；α 取”每步有效延迟”，含 kernel 内同步与协议握手）：

这解释了几个常见的观察：为什么 8 卡节点内要到 64～128 MB 才接近平台；为什么 64 卡跨机的曲线在 256 MB 以下都”看起来没跑满”；为什么 DDP 默认 25 MB 的 bucket 在单机 8 卡上刚过拐点、在 64 卡上还在爬坡（这是为什么大规模下 NCCL 会为中等消息选 Tree——把 \(2(n-1)\) 步的延迟项换成 \(2\log_2 n\) 步）。

把这些预测线叠在实测曲线上，就是第十章 sweep.sh 画的图。预测与实测的三种偏差各有含义：平台低是 β 的问题（路径、GDR、channel）；左端整体偏低（小消息 time 偏大）是 α 的问题（协议、跨 NUMA、proxy）；中段偏离是算法/协议切换的问题。下图用 ring 模型把三种偏差画在同一张对数坐标图上（8×H100 的典型参数，非实测）——α 的问题只动左半段，β 的问题只动右半段，切换点的问题只在拐点附近出现：

3. 三条参考线

下面是三种典型环境的 all_reduce busbw 曲线应有的样子。数字全部是”通常能达到”的区间，非实测；读者的机器落在区间内即为正常。

环境                          8 B time      1 MB busbw      64 MB busbw       ≥1 GB busbw (平台)     平台对应的硬件上限
────────────────────────────────────────────────────────────────────────────────────────────────────────────────────
8×H100 SXM + NVSwitch         10–20 µs      40–80 GB/s      250–400 GB/s      350–480 GB/s           NVLink 900 GB/s 双向合计
                                                                                                     = 450 GB/s 单向
8×A100 SXM + NVSwitch         12–25 µs      30–60 GB/s      150–230 GB/s      220–280 GB/s           NVLink 600 GB/s 双向合计
                                                                                                     = 300 GB/s 单向
8×GPU 仅 PCIe 4.0（无 NVLink） 20–40 µs      5–10 GB/s       10–20 GB/s        10–25 GB/s             PCIe 4.0 x16 单向 32 GB/s，
                                                                                                     多卡共享上游、经 CPU 更低
2 节点 ×8 H100，NDR 每 GPU     25–45 µs      15–25 GB/s      35–45 GB/s        40–48 GB/s             NDR 400 Gb/s = 50 GB/s 单向
2 节点 ×8 H100，HDR 每 GPU     25–45 µs      10–18 GB/s      18–23 GB/s        20–24 GB/s             HDR 200 Gb/s = 25 GB/s 单向
2 节点 ×8，每节点只有 1 张 NDR  30–50 µs      3–5 GB/s        5–6 GB/s          5–6.2 GB/s             50 GB/s 被 8 张 GPU 分摊

三点说明：

• H100 NVSwitch 上 busbw 可以超过 450 GB/s 的单向标称值。原因是 busbw 的校正系数 \(\frac{2(n-1)}{n}\) 假设的是 ring 的流量模式；开启 NVLS 后归约在 NVSwitch 内完成，步数从 \(2(n-1)\) 降到常数，完成时间比 ring 模型预估的短（每卡发 S 收 S，字节量并不比 ring 少），按 ring 系数从时间反算 busbw 就会”超标”。看到 busbw 高于单向链路上限不是测错了，而是说明 NVLS 生效了；反过来如果 H100 节点上平台只有 300 GB/s 上下，先查 NCCL_NVLS_ENABLE 与 NCCL_DEBUG=INFO 里 “NVLS multicast support is available” 那一行。
• 跨机曲线的平台由每 GPU 的网卡决定，而不是由 NVLink 决定。8 卡节点每 GPU 一张 400 Gb/s 网卡，all_reduce busbw 上限 50 GB/s；这是节点内 NVLink 平台的 1/8～1/10。多机训练的通信账几乎全部记在这一条线上。
• PCIe-only 机器的曲线不但平台低，而且形状可能不单调：跨 CPU socket 的 GPU 对经 UPI/QPI 通信，channel 数与 ring 走向对结果影响大；这类机器上 NCCL_P2P_LEVEL 与 NCCL_SHM_DISABLE 才是真正需要试的参数。

4. 从 nccl-tests 到训练

nccl-tests 给出的是”这台机器、这个 NCCL、这个消息大小”的上限。训练时实际拿到的通常更低，差距来自三处：计算 kernel 与通信 kernel 争 SM（第五篇）；框架把 buffer 切成了比拐点小的块；多个 communicator（DP、TP、PP 各一个）的流量叠加在同一条链路上。所以 nccl-tests 的正确用法是基线：它排除框架因素后告诉你硬件加 NCCL 能到哪里；训练 profiler 里的通信 kernel 时长与它比，差距就是框架侧要解释的部分。

5. 常见异常形状与原因

形状                                  最可能的原因                                    先查什么
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
大消息平台远低于参考线                 链路没走对（P2P 没走 NVLink、走了 SHM 或 CPU）    NCCL_DEBUG=INFO 里 "via P2P/…" 还是 "via SHM/…"
（节点内）                            channel 太少                                   "Connected all rings" 附近的 channel 数；TUNING 日志的 channel{Lo..Hi}
                                      NVLink 链路降级 / 掉线                          nvidia-smi nvlink -s；nvidia-smi topo -m
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
大消息平台远低于参考线                 GDR 没开（数据经 host 内存中转）                  "GPU Direct RDMA Disabled for GPU … (distance N > M)"
（跨机）                              网卡亲和错：GPU 用了跨 NUMA/跨 PCIe switch 的网卡  "via NET/IB/<dev>" 里的 dev 与 topo -m 对照
                                      NCCL_IB_HCA 没限定，混入了管理网口或不同链路类型    "NET/IB : Using [0]mlx5_…" 那一行列了哪些设备
                                      回落到 Socket                                   "Using network Socket" 而不是 "Using network IB"
                                      RoCE 没配 PFC/ECN，丢包重传                       交换机计数器、ibstat 的 port_rcv_errors
                                      链路协商在低速率                                 ibstat 的 Rate 字段
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
小消息 time 明显偏高                   协议选错（LL/LL128 被禁用，小消息走了 Simple）     TUNING 日志的 proto；NCCL_PROTO 是否被设置
（左端整体下移）                       跨 NUMA：进程绑核与 GPU/网卡不在同一 NUMA          numactl -H；进程的 CPU 亲和
                                      proxy 线程被抢占（CPU 被计算线程占满）             top 看 NCCL proxy 线程；预留 CPU 核
                                      框架没合并小消息（不是 nccl-tests 的问题）         训练侧 profiler 看通信 kernel 的大小分布
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
中段有凹陷或台阶                       算法/协议切换点选得不好                           NCCL_DEBUG_SUBSYS=TUNING 看切换发生在哪个大小
                                      channel 数随消息大小变化                          同上，看 channel{Lo..Hi}
                                      不是问题：Tree→Ring、LL128→Simple 的切换本身      切换处小幅不平滑是正常的
                                      就有几个百分点的不连续
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
多机远差于 (单机 / 节点数)             以上跨机各项之外：每节点网卡数 < GPU 数              数一下 "NET/IB : Using" 里的设备数
                                      NCCL_CROSS_NIC 配置与拓扑不匹配                    换 0/1/2 试
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
曲线抖动、in-place 与 out-of-place 差很多  其他任务干扰；GPU 降频；buffer 未对齐         nvidia-smi -q -d CLOCK；-u 参数是否设置

四、调优参数：按作用层分类

1. 参数怎么被读取

NCCL 的环境变量绝大多数经 NCCL_PARAM 宏定义（src/include/param.h）：

// src/include/param.h（2.28.9）
#define NCCL_PARAM(name, env, deftVal) \
  int64_t ncclParam##name() { \
    constexpr int64_t uninitialized = INT64_MIN; \
    static int64_t cache = uninitialized; \
    if (__builtin_expect(__atomic_load_n(&cache, __ATOMIC_RELAXED) == uninitialized, false)) { \
      ncclLoadParam("NCCL_" env, deftVal, uninitialized, &cache); \
    } \
    return cache; \
  }

三个后果值得记住。第一，值是整数（strtoll，支持 0x 前缀），且只读一次并缓存——进程启动后改环境变量无效，同一进程内两个 communicator 不可能用不同的 NCCL_BUFFSIZE。第二，ncclLoadParam（src/misc/param.cc）在 NCCL_DEBUG=INFO 且子系统含 ENV 时打印 "%s set by environment to %lld."，所以日志开头那一串 “set by environment” 就是本次生效的全部整数参数，排障时先核对这一段。第三，读取经过 ncclGetEnv，它先调用 initEnv（src/misc/param.cc）把 NCCL_CONF_FILE 指定的文件（未设时依次读 ~/.nccl.conf、/etc/nccl.conf）里的 KEY=VALUE 用 setenv(..., 0) 注入——不覆盖已有环境变量，所以优先级是 环境变量 > NCCL_CONF_FILE > ~/.nccl.conf > /etc/nccl.conf > 代码默认值。2.28 还多了一层 NCCL_ENV_PLUGIN（src/plugin/env.cc），允许外部插件接管所有环境变量的读取，集群可以用它做统一配置下发。

字符串类参数不走 NCCL_PARAM，而是各自 ncclGetEnv：NCCL_ALGO、NCCL_PROTO（src/graph/tuning.cc）、NCCL_IB_HCA（src/transport/net_ib.cc）、NCCL_SOCKET_IFNAME（src/misc/socket.cc）、NCCL_NET、NCCL_TOPO_DUMP_FILE（src/init.cc）、NCCL_GRAPH_DUMP_FILE（src/graph/search.cc）、NCCL_TOPO_FILE（src/graph/topo.cc）。NCCL_P2P_LEVEL、NCCL_P2P_DISABLE、NCCL_NET_GDR_LEVEL 三个走 src/graph/paths.cc 的 ncclGetLevel，既接受 LOC/NVL/PIX/PXB/PHB/SYS 这类路径类型名，也兼容旧式数字。

2. 算法与协议：NCCL_ALGO、NCCL_PROTO

2.28.9 的算法名（src/init.cc 的 ncclAlgoStr）：Tree、Ring、CollNetDirect、CollNetChain、NVLS、NVLSTree、PAT；协议名（ncclProtoStr）：LL、LL128、Simple。src/graph/tuning.cc 的 parseList 支持按原语指定与 ^ 排除：

NCCL_ALGO=Ring                          所有原语只用 Ring
NCCL_ALGO="ring;allreduce:tree"         默认 Ring，all_reduce 用 Tree（分号分隔按原语覆盖）
NCCL_PROTO=^LL128                       禁用 LL128（排查 LL128 在某些 PCIe 拓扑上的正确性问题时用）
NCCL_PROTO="LL,Simple;allreduce:^LL"    全局 LL+Simple，但 all_reduce 不用 LL

作用层是第四篇讲的 tuning 决策：NCCL 对每个（原语、算法、协议）组合估算时间，取最小者；设置这两个变量是把某些组合从候选集里删掉。它们是排障工具，不是调优工具：怀疑 Tree 在某个消息大小上选错了，用 NCCL_ALGO=Ring 跑一遍对照就能确认；确认之后正确的动作是把结论反馈到 tuner plugin（NCCL_TUNER_PLUGIN）或升级 NCCL，而不是在生产环境固定一个算法——固定算法会让另一段消息大小变差。

3. 并行度：NCCL_MIN_NCHANNELS、NCCL_MAX_NCHANNELS、NCCL_NTHREADS

NCCL_MIN_NCHANNELS、NCCL_MAX_NCHANNELS（src/graph/connect.cc）限定 channel 数的范围，上限是 MAXCHANNELS = 64（src/include/device.h）。每个 channel 是一个独立的 ring/tree 加一组 CUDA block 加自己的 buffer；channel 多，带宽高（更多 SM、更多并发的 NIC QP），但小消息延迟略增、占用的 SM 更多、与计算 kernel 的争抢更明显。第五篇讨论重叠时讲过：训练里通信 kernel 多占一个 SM 就少一个 SM 给 GEMM。

NCCL_NTHREADS（src/graph/tuning.cc）是每个 channel 的 CUDA block 线程数，Simple 协议默认 512（NCCL_SIMPLE_MAX_NTHREADS），PCIe 带宽受限的拓扑上默认降为 256；LL128 上限 640。它必须是 32 的倍数且在 64～640 之间，几乎没有理由手动设。

这三者是”两种账”冲突最直接的例子：带宽的账要 channel 多，延迟的账和计算重叠的账要 channel 少。nccl-tests 的大消息平台上不去时可以试 NCCL_MAX_NCHANNELS=32 看看是否有提升，有则说明拓扑探测给的默认 channel 数偏保守；但在训练里应该用 profiler 验证通信与计算的重叠没有变差再决定。

4. 缓冲：NCCL_BUFFSIZE

NCCL_BUFFSIZE（src/init.cc）是每个 channel 每个方向的 Simple 协议 buffer，默认 DEFAULT_BUFFSIZE = 4 MiB，分成 NCCL_STEPS 个 slot 做流水。它决定了跨机传输的每一步在途字节数：跨机 α ≈ 10 µs、β = 50 GB/s 时，一个 4 MiB 的 buffer 分 8 步、每步 512 KB，在途数据 512 KB / 50 GB/s ≈ 10 µs，刚好覆盖一步的延迟。长距离（跨机房、高延迟 RoCE）链路上带宽平台上不去、且 ib_write_bw 正常，可以试 8 MiB 或 16 MiB；节点内 NVLink 上改它几乎没有效果。NCCL_LL_BUFFSIZE、NCCL_LL128_BUFFSIZE 同理但更少需要动。

5. 路径选择：NCCL_P2P_LEVEL、NCCL_P2P_DISABLE、NCCL_SHM_DISABLE、NCCL_NET_GDR_LEVEL

这组变量作用在第四篇讲的 transport 选择上，取值是第二篇讲的 nvidia-smi topo -m 里那套路径类型（src/graph/topo.cc 的 topoPathTypeStr：LOC NVL NVB C2C PIX PXB P2C PXN PHB SYS NET DIS）：

• NCCL_P2P_LEVEL：两张 GPU 距离不超过这个级别才用 CUDA P2P（NVLink 或 PCIe 直通），否则走 SHM。默认 PXB（src/graph/paths.cc 的 p2pLevel = PATH_PXB），即同一 PCIe switch 下允许，跨 CPU（PHB、SYS）不允许——因为跨 CPU 的 P2P 在很多平台上比经 host 内存中转还慢。NCCL_P2P_DISABLE=1 等价于 NCCL_P2P_LEVEL=LOC。
• NCCL_SHM_DISABLE=1（src/transport/shm.cc）禁用共享内存传输，两张不能 P2P 的 GPU 会被迫走网络（哪怕在同一台机器上）。只在怀疑 SHM 路径有 bug 时用。
• NCCL_NET_GDR_LEVEL：GPU 与网卡距离不超过这个级别才启用 GPUDirect RDMA，默认也是 PXB（netGdrLevel = PATH_PXB）。GPU 和网卡跨了 CPU 时默认关 GDR，日志里出现 GPU Direct RDMA Disabled for GPU x / HCA y (distance N > M)，数据经 host 内存中转，跨机 busbw 掉到 PCIe/内存拷贝的水平。正确的修法是让 GPU 用对的网卡（拓扑亲和、NCCL_IB_HCA），而不是把 level 放宽到 SYS——放宽后 GDR 跨 CPU 通常比不开还慢。

6. 网络：NCCL_IB_HCA、NCCL_IB_GID_INDEX、NCCL_IB_TC、NCCL_IB_QPS_PER_CONNECTION、NCCL_IB_SPLIT_DATA_ON_QPS

全部在 src/transport/net_ib.cc：

• NCCL_IB_HCA：限定 NCCL 使用的 IB 设备与端口，逗号分隔，mlx5_0:1 形式指定端口；前缀 ^ 排除、= 精确匹配（否则是前缀匹配，mlx5_1 会同时匹配 mlx5_1 和 mlx5_10）。生产环境应当显式设置：不设时 NCCL 会枚举所有 ibv_get_device_list 返回的设备，包括存储网、管理网上的 RoCE 网卡，出现 Attempted to merge incompatible devices … Try selecting NICs of only one link type using NCCL_IB_HCA 的告警或直接把流量分到慢网卡上。
• NCCL_IB_GID_INDEX：RoCE v2 必须指定正确的 GID index（对应 RoCE v2 + IPv4 的那一项，show_gids 可查）；IB 上不需要。设错的症状是 QP 建立失败或跨机完全不通。2.28 里默认 -1 表示自动选择，多数 rdma-core 版本能选对，但混合 v1/v2 的环境仍需手动。
• NCCL_IB_TC：Traffic Class，RoCE 上用于把 NCCL 流量映射到配置了 PFC/ECN 的优先级队列（常见值 106 对应 DSCP 26），必须与交换机配置一致；IB 上一般不用。NCCL_IB_SL 是 IB 上的对应物。
• NCCL_IB_QPS_PER_CONNECTION（默认 1）：每个连接开几条 QP。多条 QP 在有多路径（ECMP）的以太网 fabric 上可以把一条流打散到多条链路，缓解哈希冲突；NCCL_IB_SPLIT_DATA_ON_QPS（默认 0）决定是把每个消息切开分到各 QP（1）还是各消息轮转（0）。IB fabric 有自适应路由时不需要；RoCE 上遇到多机带宽不稳、且交换机侧看到链路利用不均时试 2 或 4。
• 同文件里还有 NCCL_IB_TIMEOUT（默认 20）、NCCL_IB_RETRY_CNT（默认 7）、NCCL_IB_AR_THRESHOLD、NCCL_IB_PCI_RELAXED_ORDERING、NCCL_IB_ROCE_VERSION_NUM、NCCL_IB_MERGE_NICS；除 timeout 在大规模 RoCE 上偶尔需要调大之外，其余按默认。

7. Socket 与 bootstrap：NCCL_SOCKET_IFNAME、NCCL_SOCKET_NTHREADS、NCCL_NSOCKS_PERTHREAD

NCCL_SOCKET_IFNAME（src/misc/socket.cc 的 ncclFindInterfaces）指定 bootstrap（交换 ncclUniqueId 之后建立 out-of-band 连接）以及 Socket transport 用的网口，支持前缀匹配与 ^ 排除。设错的典型症状是初始化 hang 在 bootstrap 阶段（NCCL_DEBUG=INFO 停在 Bootstrap: Using <ifname> 之后不动）或者选中了 docker0 之类不可路由的接口。它不影响 IB 数据路径的选择——那是 NCCL_IB_HCA 的事。

NCCL_SOCKET_NTHREADS、NCCL_NSOCKS_PERTHREAD（src/transport/net_socket.cc）只在数据走 Socket transport 时有意义：多线程多 socket 才能把 TCP 推到接近线速。如果你在 IB/RoCE 集群上发现自己需要调它们，说明数据根本没走 RDMA，先回去看 “Using network” 那一行。

8. 特殊功能：NCCL_CROSS_NIC、NCCL_NVLS_ENABLE、NCCL_COLLNET_ENABLE

• NCCL_CROSS_NIC（src/graph/search.cc，默认 2）：ring 跨机时是否允许从一张网卡出、另一张网卡进。0 强制同网卡（rail-optimized 网络，即每张网卡接到独立的 leaf 交换机时应设 0）；1 允许；2 让 NCCL 按拓扑决定。设错在 rail 网络上会让流量跨 spine，多机 busbw 下降且不稳定。
• NCCL_NVLS_ENABLE（src/transport/nvls.cc，默认 2 = 自动）：Hopper 及之后 + NVSwitch 上启用 NVLink SHARP（多播归约）。如果第三节看到 H100 平台明显低于 400 GB/s，日志里没有 NVLS multicast support is available，检查驱动版本与 Fabric Manager。
• NCCL_COLLNET_ENABLE（src/init.cc）：启用 SHARP（IB 交换机内归约）之类的 CollNet 插件，需要网络侧支持与 NCCL_NET_PLUGIN；一般集群不具备，保持默认。

9. 该动、少动、几乎不动

类别            参数                                              理由
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
生产环境应显式   NCCL_IB_HCA                                       不设会枚举所有 RDMA 设备，混入慢网卡
设置的           NCCL_SOCKET_IFNAME                                bootstrap 选错网口是初始化 hang 的第一原因
                NCCL_IB_GID_INDEX（RoCE v2）                       自动选择不总是对
                NCCL_IB_TC / NCCL_IB_SL（RoCE 有 PFC/ECN 时）        流量必须落在配了无损的队列里
                NCCL_DEBUG=WARN（生产）/ INFO（排障）               见第五章
                NCCL_DEBUG_FILE（多机排障时）                        不然 64 个进程的日志混在一起
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
排障时用来对照   NCCL_ALGO / NCCL_PROTO                            确认某个算法/协议是否是问题源；确认后不要留在生产配置里
的（跑完就撤）   NCCL_P2P_DISABLE / NCCL_SHM_DISABLE / NCCL_P2P_LEVEL  确认路径选择是否是问题源
                NCCL_NET_GDR_LEVEL                                确认 GDR 是否是问题源；修法是改亲和不是放宽 level
                NCCL_MIN/MAX_NCHANNELS                            确认 channel 数是否偏保守；训练里要验证重叠
                NCCL_CROSS_NIC                                    rail 网络上试 0
                NCCL_IB_QPS_PER_CONNECTION                        RoCE ECMP 不均时试 2/4
──────────────────────────────────────────────────────────────────────────────────────────────────────────────
几乎永远不该动   NCCL_NTHREADS / NCCL_LL128_NTHREADS               默认值是按拓扑算的
                NCCL_BUFFSIZE / NCCL_LL_BUFFSIZE                  只有高延迟长链路才需要
                NCCL_IB_TIMEOUT / NCCL_IB_RETRY_CNT               调大只是把"报错"变成"更久的 hang"
                NCCL_NVLS_ENABLE / NCCL_COLLNET_ENABLE            让 NCCL 自动判断；设 0 只用于验证它是不是 bug 源
                NCCL_TOPO_FILE                                    手写拓扑文件只用于虚拟机里探测不到拓扑的情形

一条总原则：每个改动都要用 nccl-tests 的前后两条曲线证明，并且要看整条曲线而不是一个点。把 NCCL_ALGO=Ring 写进生产配置的团队，通常是三年前在某台机器上的某个消息大小上看到过 5% 的提升。

五、日志：NCCL_DEBUG 的三层

1. 级别

src/debug.cc 的 ncclDebugInit 解析 NCCL_DEBUG，接受（不区分大小写）VERSION、WARN、INFO、ABORT、TRACE：

• 未设置：什么都不打，连版本都不打。
• VERSION：只打一行 NCCL version 2.28.9+cuda12.x。
• WARN：只打 WARN 级别——错误与异常路径。生产默认应当是 WARN：开销可忽略，而 hang 或性能问题时那一行 NCCL WARN 往往就是答案（Cuda failure、Got completion from peer … with error 12、Attempted to merge incompatible devices、socket … Connection reset）。
• INFO：加上初始化阶段的全部决策——拓扑、ring/tree、transport、算法协议表、环境变量。这是排性能问题的级别。它在初始化时打几百到几千行，稳态时几乎不打（除非 NCCL_DEBUG_SUBSYS 含 COLL），所以排障期开着它是可以接受的。
• ABORT：在第一个 WARN 处 abort 进程，得到 core dump。
• TRACE：每一次 API 调用与内部步骤。注意 src/include/debug.h 里 TRACE(...) 宏只在编译时带 -DENABLE_TRACE（make TRACE=1）才有实体，发行版二进制里 NCCL_DEBUG=TRACE 只多出 TRACE_CALL 那一类 NCCL CALL ncclAllReduce(...) 的 API 调用记录，其余与 INFO 相同。

NCCL_WARN_ENABLE_DEBUG_INFO=1 是一个有用的组合：平时按 WARN 运行，一旦出现 WARN 就把级别提到 INFO，之后的日志会带上更多上下文。

每行日志的格式（ncclDebugLog）：<hostname>:<pid>:<tid> [<cudaDev>] NCCL INFO ...，WARN 行多出 <file>:<line>。多机排障时用 <hostname>:<pid> 定位到 rank，用 [<cudaDev>] 定位到卡。

2. 子系统：NCCL_DEBUG_SUBSYS

INFO 默认只开 INIT、BOOTSTRAP、ENV 三个子系统（ncclDebugInit 里的 tempNcclDebugMask）。NCCL_DEBUG_SUBSYS 用逗号分隔选择，^ 前缀取反。2.28.9 的完整列表：

INIT       初始化流程、communicator 创建、"Connected all rings/trees"
COLL       每一次集合通信调用：opCount、buffer、count、datatype、op、comm、stream（稳态大量输出）
P2P        P2P transport 的连接建立（"via P2P/CUMEM"）
SHM        SHM transport
NET        网络 transport：设备枚举、"via NET/IB/…/GDRDMA"、GDR 开关
GRAPH      拓扑图与搜索：ring/tree 的构造，"Ring 00 : 3 -> 0 -> 1"
TUNING     算法/协议表；enqueue 时 rank 0 打 "AllReduce: N Bytes -> Algo Ring proto LL128 channel{Lo..Hi}"
ENV        环境变量读取："NCCL_XXX set by environment to N"
ALLOC      内存分配
CALL       API 调用（TRACE 级别）
PROXY      proxy 线程
NVLS       NVLink SHARP
BOOTSTRAP  bootstrap 网络
REG        buffer 注册
PROFILE    profiler 插件
RAS        RAS 子系统
ALL        全部

三个常用组合：

• 排性能：NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,GRAPH,NET,TUNING——拓扑决策、transport、每个消息大小选了什么。
• 找”是哪一次 all_reduce”：NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=COLL——src/enqueue.cc 会为每次调用打 AllReduce: opCount 3f2 sendbuff … count 6553600 datatype 7 op 0 root 0 comm 0x… [nranks=64] stream 0x…。opCount 是每个 communicator 内单调递增的序号；hang 之后比较各 rank 最后一行的 opCount，就能看出谁少了一次。代价是每次调用一行日志，训练时每步几百行，只在复现问题时开。
• 排网络：NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=NET,ENV。

3. 按 rank 分文件：NCCL_DEBUG_FILE

NCCL_DEBUG_FILE 指定输出文件，ncclDebugInit 里展开两个占位符：%h 主机名、%p 进程号（%% 转义为 %）。多机多进程时必须用它，否则 stdout 里 64 个进程的行交织在一起无法阅读：

NCCL_DEBUG=INFO NCCL_DEBUG_FILE=/shared/logs/nccl.%h.%p.log

文件用 setlinebuf 行缓冲，进程崩溃时不会丢已写的行。只有 NCCL_DEBUG 高于 VERSION 时才会创建文件。注意没有 %r（rank）占位符——NCCL 在读这个变量时还不知道自己的 rank；要按 rank 命名，用启动器的环境变量（torchrun 的 RANK、mpirun 的 OMPI_COMM_WORLD_RANK）拼路径，例如在 wrapper 脚本里 export NCCL_DEBUG_FILE=/shared/logs/nccl.rank${RANK}.log。

4. 一份 INFO 日志里要找的八行

拿到一份 NCCL_DEBUG=INFO 日志，不要从头读，直接 grep 这八样东西（第四篇的 nccl_log_reader.py 做的就是这个）：

1  NCCL version 2.28.9+cuda12.x                             版本对不对；多机各节点是否一致
2  NCCL_XXX set by environment to …                          生效的参数就这些，别的都是默认
3  Bootstrap: Using eth0:10.0.0.1<0>                         bootstrap 网口
4  NET/IB : Using [0]mlx5_0:1/IB [1]mlx5_1:1/IB … ; OOB eth0  数据网卡列表：数量、链路类型（IB/RoCE）、是否混入了不该有的
   Using network IB                                          还是 "Using network Socket"？
5  GPU Direct RDMA Enabled/Disabled for GPU x / HCA y        GDR 是否生效，distance 多少
6  Channel 00/0 : 0[0] -> 1[1] via P2P/CUMEM                 节点内路径：P2P 还是 SHM
   Channel 00/0 : 7[7] -> 8[0] [send] via NET/IB/0/GDRDMA    跨机路径：哪张网卡、带不带 GDRDMA
7  Connected all rings, use ring PXN 0 GDR 1                  ring 建好；channel 数看 "Channel NN/…" 的最大编号
   Connected all trees
8  （SUBSYS 含 TUNING 时）AllReduce: 33554432 Bytes -> Algo Ring proto Simple channel{Lo..Hi}={0..15}
                                                             这个大小选了什么；与曲线上的台阶对照

前七行都在初始化阶段，第八行在每次调用（只有 rank 0 打）。曲线异常时，几乎总能在这八行里找到与预期不符的一处。

5. NCCL 自己的运行时视角：RAS 与 ncclras

NCCL 2.24 起内置了 RAS（Reliability, Availability, Serviceability）子系统，源码在 src/ras/，默认启用（src/bootstrap.cc 的 NCCL_PARAM(RasEnable, "RAS_ENABLE", 1)）。每个进程里一个 RAS 线程，进程之间用 TCP 组成一张监控网，与 NCCL 的数据路径完全独立。随 NCCL 一起编译的命令行工具 ncclras（src/ras/client.cc，make install 装到 bin/）连接到任一进程的 RAS 端口（默认 28028，NCCL_RAS_ADDR 可改），输出全作业的状态：

ncclras -v          # 或 ncclras -f json
Job summary
===========
  Nodes  Processes         GPUs  Processes     GPUs
(total)   per node  per process    (total)  (total)
      8          8            1         64       64
Communicators
...

src/ras/client_support.cc 为每个 communicator 汇总各 rank 的 collOpCounts（按原语类型计数的集合通信次数）并比较，不一致时标记 MISMATCH，有 rank 没响应时标记 INCOMPLETE，并把与多数不同的少数 rank 作为 outlier 列出。这与下一章的 Flight Recorder 在原理上相同——比较各 rank 的操作计数找掉队者——但它不依赖 PyTorch，对 nccl-tests、vLLM 的 PyNccl、自写的 NCCL 程序都有效，而且是实时查询而不是 timeout 后 dump。它给不出 Python 栈和 tensor 形状，所以在 PyTorch 训练里两者是互补的。

六、hang 的分类与定位工具

1. 为什么 hang 是集合通信的默认失败模式

第四篇讲过 NCCL kernel 的工作方式：每个 channel 的 block 在 device 内存的 flag 上自旋，等对端把数据和 flag 写过来；数据到了就处理、再向下一跳写。没有任何一处有超时——kernel 不知道”对端应该在多久内到”，也没有心跳。第五篇讲过 CPU 侧：ncclAllReduce 在 enqueue 之后立即返回，Python 线程继续往下走，直到下一个需要 GPU 结果的同步点（.item()、torch.cuda.synchronize()、下一步的 loss.backward() 隐含的依赖）才阻塞在 cudaStreamSynchronize 或 event 等待上。

于是一个 rank 少调一次 all_reduce 的后果是：其他 63 个 rank 的 kernel 在等它的数据，它们的 CPU 线程阻塞在同步点，日志停在同一行；而它自己已经跑到下一次 all_reduce 并阻塞在那里——它的 kernel 在等其他人。所有人都在等，没有人会报错。唯一会打破这个僵局的是 ProcessGroupNCCL 的 watchdog：每个 WorkNCCL 从创建起计时，超过 opTimeout_ 就判定 timeout（WorkNCCL::checkTimeout），日志里出现 Watchdog caught collective operation timeout: WorkNCCL(SeqNum=…, OpType=ALLREDUCE, NumelIn=…, NumelOut=…, Timeout(ms)=600000) ran for 600xxx milliseconds before timing out.——这条会在每一个 rank 上出现，因为每个 rank 都有一个等不到的 work。

所以 hang 的排查等价于回答一个问题：是所有 rank 都在等同一个不存在的东西（网络/硬件/死进程），还是有一个 rank 在等别的东西（调用不一致）？ 前者看基础设施，后者看代码。

2. 六类 hang

类别                          机制                                                   现象特征
─────────────────────────────────────────────────────────────────────────────────────────────────────────────────
A 集合通信参数不一致           rank 3 传 1 MB+4 KB，其他传 1 MB；NCCL 不校验 count      要么 hang（多数），要么静默完成给出错误结果（少数，
  （tensor 大小 / dtype 不同）  跨 rank 是否一致，ring 上各 rank 按自己的 count 切块，   取决于协议与切块对齐）；Flight Recorder 报
                              收发字节数对不上                                        SIZE_OR_SYNTAX_MISMATCH 或 COLLECTIVE_DTYPE_MISMATCH
─────────────────────────────────────────────────────────────────────────────────────────────────────────────────
B 调用顺序 / 次数不一致        rank 2 因 if 分支跳过一次 all_reduce（或多做一次）；     各 rank 停在不同的序号上；rank 2 的 collective_seq_id
                              之后它的第 k 次配的是别人的第 k+1 次                    比别人少 1；若各次大小恰好相同则要到最后一次才 hang
─────────────────────────────────────────────────────────────────────────────────────────────────────────────────
C send/recv 不配对             rank 0 send 给 1，rank 1 没有 recv（或 recv 自 2）；     只有涉及的 rank 先停；其他 rank 走到下一次集合通信
                              流水线并行的 stage 边界最常见                           后被连锁拖住；Flight Recorder 里 is_p2p 条目找不到配对
─────────────────────────────────────────────────────────────────────────────────────────────────────────────────
D 两个 communicator 交叉等待   rank A 先在 comm1 上 all_reduce 再在 comm2 上；rank B    经典死锁；各 rank 各自"合法"，只是顺序不同；
                              顺序反过来；两个 kernel 各占着 SM 等对方                 py-spy 看到不同 rank 停在不同 PG 的调用上
─────────────────────────────────────────────────────────────────────────────────────────────────────────────────
E 一个 rank 崩溃 / 被 kill /   rank 17 OOM 退出、被 OOM killer 杀、Python 异常退出但    ps 看进程数少了一个；dmesg 有 Killed process；
  卡在非通信代码               没触发 abort；或者活着但卡在 dataloader / checkpoint /   py-spy 看它的栈不在 c10d 里；其他所有 rank 停在同一序号
                              GIL 死锁 / 第三方库的 mutex                              且状态一致
─────────────────────────────────────────────────────────────────────────────────────────────────────────────────
F 网络或硬件真的断了           链路 down、网卡 reset、交换机重启、GPU Xid 错误、        NCCL WARN：Got completion from peer … with error /
                              NVLink 错误                                            Cuda failure / socket closed by remote；dmesg 有 mlx5 / Xid；
                                                                                    ibstat 的 State 不是 Active；所有 rank 停在同一序号

A～D 是代码问题，出现在特定步数、可复现，且各 rank 的”最后一次操作”不一致；E、F 是环境问题，随机出现，各 rank 的”最后一次操作”一致（除了崩掉的那个没有）。这个区分是第九章决策树的第一层。

3. 看栈：py-spy、gdb、cuda-gdb

hang 发生后进程都还在（watchdog 默认 10 分钟后才动手），先看栈。三个工具看三层：

# Python 栈：每个 rank 停在哪一行 Python 代码。--native 会带上 C 栈
py-spy dump --pid <pid>                    # 不需要 root，不 attach，秒级
py-spy dump --pid <pid> --native | head -80

# C++ 栈：确认是阻塞在 cudaStreamSynchronize / ncclCommAbort / TCPStore wait 还是别的 mutex
gdb -p <pid> -batch -ex "thread apply all bt" 2>/dev/null > gdb.rank${RANK}.txt

# kernel 是否在跑：NCCL kernel 在自旋时 nvidia-smi 的 GPU-Util 是 100%、SM 有占用但没有进展
nvidia-smi --query-gpu=utilization.gpu,memory.used --format=csv
cuda-gdb -p <pid>   →  (cuda-gdb) info cuda kernels     # 看到 ncclDevKernel_* 一直在，即 kernel 在等

把 64 个 rank 的 py-spy dump 收集到一处（用 pdsh / clush / Slurm 的 srun --overlap），然后比较：

• 63 个停在 train.py:312 loss.backward() → .../c10d ... work.wait()，1 个停在 dataloader.py 的 queue.get() 或 torch.save：类 E，这个 rank 就是答案；
• 所有 64 个停在同一行同一个 all_reduce：类 E（崩了的那个不在列表里，数进程数）或类 F（看 dmesg、ibstat、NCCL WARN）；
• 一部分停在 all_reduce、一部分停在 send/recv：类 C；
• 停在不同 process group 的调用上：类 D；
• 都停在 all_reduce 但行号不同（一个在 step k 的 forward、其他在 step k 的 backward）：类 A 或 B，需要 Flight Recorder 确认是哪一次。

py-spy 的局限是只能看到”卡在哪个调用”，看不到”卡的这个调用和别人的哪个调用配对”，也看不到已经完成的历史。这正是 Flight Recorder 补上的部分。

4. Flight Recorder：记录什么、怎么开、怎么 dump

Flight Recorder 是 ProcessGroupNCCL 里的一个环形缓冲（torch/csrc/distributed/c10d/FlightRecorder.hpp 的 FlightRecorder<EventType>，CUDA 实例在 FlightRecorderCuda.cpp）。每次 initWork 创建 WorkNCCL 时调用 FlightRecorderCUDA::get()->recordWithResetEnabled(...) 写入一条 Entry，字段包括：

pg_id / pg_name            哪个 process group（名字与 desc）
collective_seq_id          该 PG 内集合通信的序号（每次 kernel launch 加 1；coalesced 组算一次）
p2p_seq_id                 该 PG 内 send/recv 的序号
op_id                      逻辑操作序号（coalesced 组内每个 op 一个）
profiling_name             "nccl:all_reduce" 这类名字
input_sizes / dtypes       输入 tensor 的形状与类型
output_sizes / dtypes      输出
traceback_                 调用时的 Python 栈（CapturedTraceback::gather；TORCH_NCCL_TRACE_CPP_STACK=1 加 C++ 栈）
time_created_ns            enqueue 时间
timeout_ms                 这条 work 的 timeout
start_ / end_ event        用于查询状态：scheduled / started / completed
time_discovered_started    watchdog 发现 kernel 已开始的时间
time_discovered_completed  watchdog 发现 kernel 已完成的时间
retired_                   已从 workMetaList_ 移出；retired 但未 completed = 超时了
thread_id / thread_name    发起调用的线程

关键的一点：它记录的是每一次调用，不只是失败的那一次。所以 dump 出来的是”最近 N 次集合通信的完整账本”，包括每次的形状、每次的 Python 栈、每次是否完成。把 64 份账本按 collective_seq_id 对齐，就能看到在哪个序号上谁的记录与别人不同。

开关与默认值（PyTorch 2.12 的 ProcessGroupNCCL.cpp、FlightRecorder.hpp 与 FlightRecorder.cpp；其他版本不同，注意核对）：

TORCH_NCCL_TRACE_BUFFER_SIZE     环形缓冲条数；2.12 默认 2000（即默认开启），设 0 关闭
                                 （别名 TORCH_FR_BUFFER_SIZE，优先于旧名）
TORCH_NCCL_DUMP_ON_TIMEOUT       watchdog 判定 timeout 或异常时 dump；2.12 默认 true
TORCH_NCCL_ENABLE_MONITORING     heartbeat monitor 线程，负责执行 dump 与在 watchdog 卡死时杀进程；2.12 默认 true
TORCH_NCCL_DEBUG_INFO_TEMP_FILE  dump 文件前缀，后接全局 rank 号；2.12 默认 $XDG_CACHE_HOME/torch/comm_lib_trace_rank_
                                 （XDG_CACHE_HOME 未设时为 ~/.cache/torch/comm_lib_trace_rank_）；别名 TORCH_FR_DUMP_TEMP_FILE
TORCH_NCCL_TRACE_CPP_STACK       同时记录 C++ 栈（有开销），默认 false（别名 TORCH_FR_CPP_STACK）
TORCH_NCCL_HEARTBEAT_TIMEOUT_SEC monitor 判定 watchdog 自己卡死的时间，默认 480 秒
TORCH_NCCL_WAIT_TIMEOUT_DUMP_MILSEC  dump 的时间预算，默认 15000 ms
TORCH_NCCL_DEBUG_INFO_PIPE_FILE  设置后 monitor 在 <前缀><全局rank>.pipe 上监听一个命名管道，往里写就触发 dump——不用等 timeout

生产训练建议显式写上这四个，不依赖默认值随版本变动：

export TORCH_NCCL_TRACE_BUFFER_SIZE=2000
export TORCH_NCCL_DUMP_ON_TIMEOUT=1
export TORCH_NCCL_ENABLE_MONITORING=1
export TORCH_NCCL_DEBUG_INFO_TEMP_FILE=/shared/fr_dumps/${JOB_ID}/trace_   # 共享存储，后接 rank 号

dump 的触发流程（ProcessGroupNCCL.cpp 的 Watchdog::runLoop 与 HeartbeatMonitor::runLoop）：

watchdog 发现某个 work 超时
  → 打 "Watchdog caught collective operation timeout … " 与 "failure detected by watchdog at work sequence id: N
     PG status: last enqueued work: N, last completed work: K"
  → work.printTraceback() 打出这次调用的 Python 栈
  → broadcastDumpSignal(): 在全局 TCPStore 上写 key "exception_dump"（kStoreDumpKey），通知所有 rank
  → 本地 shouldDump_ = true；sleep 4 × TORCH_NCCL_WAIT_TIMEOUT_DUMP_MILSEC（默认 60 秒）给 dump 留时间
  → 之后按 TORCH_NCCL_ASYNC_ERROR_HANDLING（默认 3 = SkipCleanUp）抛异常，进程退出
每个 rank 的 monitor 线程每 TORCH_NCCL_COORD_CHECK_MILSEC（1 秒）检查 TCPStore 上的 "exception_dump"
  → 看到就调用 dumpDebuggingInfo()，把 FlightRecorder 的内容 pickle 后交给 DebugInfoWriter::write 写文件

这条流程里有两个线程、两类 rank 在协作——超时的 rank 的 watchdog 只负责”喊”，真正写文件的是每个 rank 自己的 monitor 线程，包括那个没有超时的掉队者：

%% 图：Flight Recorder 的 dump 流程：超时 rank 的 watchdog 只负责经 TCPStore 广播信号，每个 rank 的 monitor 线程各自写文件
sequenceDiagram
    participant WD as rank k 的 watchdog
    participant ST as TCPStore
    participant MK as rank k 的 monitor
    participant M5 as rank 5 的 monitor (掉队者)
    Note over WD: checkTimeout() 为真：某 work 超过 opTimeout_
    WD->>WD: 打 timeout 日志, printTraceback()
    WD->>ST: broadcastDumpSignal() 写 key exception_dump
    WD->>MK: shouldDump_ = true
    loop 每 TORCH_NCCL_COORD_CHECK_MILSEC (1 s)
        MK->>ST: 查 exception_dump
        M5->>ST: 查 exception_dump
    end
    ST-->>MK: 有信号
    ST-->>M5: 有信号
    MK->>MK: dumpDebuggingInfo() 写 trace_k
    M5->>M5: dumpDebuggingInfo() 写 trace_5 (没超时也 dump)
    Note over WD,MK: watchdog sleep 4 x WAIT_TIMEOUT_DUMP_MILSEC (60 s) 给 dump 留时间
    WD->>WD: 抛 DistBackendError (ASYNC_ERROR_HANDLING=3), 进程退出

于是所有 rank——包括那个没超时、只是掉队的 rank——都会 dump。这是它能对齐的前提。如果 TCPStore 已经不可用（rank 0 死了），其他 rank 收不到信号，就只能靠各自的 timeout 触发本地 dump。

不想等 10 分钟 timeout 时，有两条手动触发的路：torch._C._distributed_c10d._dump_nccl_trace()（Python 侧，返回 pickle bytes，需要 Python 线程能响应——hang 在 .item() 时它不能）；或者设置 TORCH_NCCL_DEBUG_INFO_PIPE_FILE=/tmp/fr_pipe，hang 时 echo 1 > /tmp/fr_pipe<rank>.pipe，monitor 线程会 dump（它不依赖 Python 线程；管道只在 uid 为 0 的默认 PG 上创建，按全局 rank 命名）。后者是复现实验里最方便的。

5. fr_trace.py 怎么对齐

分析工具在 torch/distributed/flight_recorder/fr_trace.py，逻辑在 components/（loader.py 读文件、builder.py 建库、utils.py 匹配）：

python -m torch.distributed.flight_recorder.fr_trace /shared/fr_dumps/${JOB_ID}/ -p trace_
# 或 python torch/distributed/flight_recorder/fr_trace.py <dir> --prefix trace_
#   -j / --just_print_entries        只打印所有 rank 的条目，不做匹配（先用它看一眼）
#   --selected-ranks 3 17            只看这几个 rank
#   --pg-filters default tp          只看某些 PG
#   --allow-incomplete-ranks         少了几个 rank 的文件也分析
#   --print_stack_trace              打出条目的 Python 栈
#   --mismatch_cap 10                最多打印前 10 处不匹配
#   -o db.pkl                        导出扁平数据库

它的算法：把每个 rank 的条目按 PG 分组、按 collective_seq_id 排序，然后对每个 PG 从最小的序号开始，取所有 rank 在该序号上的条目互相比对（utils.py 的 match_one_event），得到 types.py 里的 MatchState：

FULLY_MATCHED               原语、形状、dtype、状态都一致
COLLECTIVE_TYPE_MISMATCH    rank 3 在这个序号上是 all_gather，其他是 all_reduce    → 类 B（顺序不一致）
SIZE_OR_SYNTAX_MISMATCH     形状不一致，或 send/recv 不配对                       → 类 A 或 C
COLLECTIVE_DTYPE_MISMATCH   dtype 不一致                                          → 类 A
COLLECTIVE_STATE_MISMATCH   有的 completed、有的 scheduled/started                 → 有 rank 根本没开始这次操作（类 E 的"卡在别处"）
UNDECIDED                   alltoall 这类需要看全部 rank 才能判断

输出里第一处非 FULLY_MATCHED 的序号和 culprit rank 就是答案。把 64 份账本按序号排成一张表，三大类 hang 各有一眼可辨的签名——这正是第九章决策树 hang 分支第一层”一致还是不一致”的判据：

fr_trace.py 的视角：一行一个 rank，一格一条 entry；格内是 input_sizes（元素数）
与状态：C = completed，S = started（kernel 已启动、在等对端）

① 类 A（代码）：同一序号上形状不同 → SIZE_OR_SYNTAX_MISMATCH，culprit = rank 3
          seq 4        seq 5        seq 6
rank 0    [1M] C       [2M] C       [3M] S
rank 1    [1M] C       [2M] C       [3M] S
rank 3    [1M] C       [2M] C       [3M+1K] S   ← 只有它不同，看这条的 Python 栈
rank 7    [1M] C       [2M] C       [3M] S
  （类 B 的样子相同：跳过一次的 rank 在该序号上顶着"下一步"的形状 [4M]）

② 类 E-卡住（环境）：rank 5 比别人少一条，且它的最后一条是 C
          seq 4        seq 5        seq 6
rank 0    [1M] C       [2M] C       [3M] S
rank 1    [1M] C       [2M] C       [3M] S
rank 5    [1M] C       [2M] C         --        ← 没调用 seq 6：卡在 torch.save
rank 7    [1M] C       [2M] C       [3M] S
  （类 E-崩溃：rank 5 连 dump 文件都没有，fr_trace 加 --allow-incomplete-ranks）

③ 类 F（网络 / 硬件）：全部 rank 完全一致，都停在 seq 6、都是 S
          seq 4        seq 5        seq 6
rank 0    [1M] C       [2M] C       [3M] S
rank 1    [1M] C       [2M] C       [3M] S
 ...       ...          ...          ...
rank 63   [1M] C       [2M] C       [3M] S      ← 无 culprit：看 WARN、dmesg

对类 B 有一个细节：如果每一步的 all_reduce 大小完全相同，跳过一次的 rank 在后续每个序号上的形状都与别人一致，匹配会一路 FULLY_MATCHED，直到最后一个序号上别人有条目、它没有——这时报的是 COLLECTIVE_STATE_MISMATCH 或缺失条目，指向同一个 rank，只是不能告诉你它是在哪一步跳过的。真实训练里 DDP bucket、FSDP 各层参数的形状各不相同，所以这个歧义很少出现；第十章 hang_lab/ 的脚本故意让每步大小不同，就是为了让匹配结果无歧义。

6. TORCH_NCCL_DESYNC_DEBUG

Flight Recorder 之前的机制，仍然可用。开启后（ProcessGroupNCCL.cpp 的 DesyncDebugger；同时会强制 enableTiming_），每个 work 开始和结束时把 (seq, 原语名) 写到 TCPStore 上的 per-rank key（TraceUtils.h 的 getTraceStartKey/getTraceEndKey）。timeout 时 retrieveDesyncReport 读所有 rank 的 key，输出：本 rank 停在哪个序号、哪些 rank 没有记录（missing）、哪些 rank 落后（lagging）、以及一张 seq → rank → start/end 的快照。它比 Flight Recorder 轻（只记最后一次，写 TCPStore）、也粗（没有形状、没有栈、只在 timeout 的那个 rank 上报告）。TCPStore 在几千 rank 时会成为热点，所以新版本更推荐 Flight Recorder；两者可以同时开。

类 E 的两个子类在这些工具下的签名不同：”崩了”的 rank 没有 dump 文件（fr_trace.py 加 --allow-incomplete-ranks），答案在它的 stderr 与 dmesg 里，TORCH_NCCL_PROPAGATE_ERROR=1 可让 watchdog 把错误经 TCPStore 广播给同 PG 的其他 rank 提前退出；”活着但卡在别处”的 rank 有 dump，但比别人少一条且最后一条是 completed，其他 rank 的最后一条是 started——这是”卡在通信之外”的签名，下一章用它解释 checkpoint 卡住为什么变成 NCCL timeout。

七、timeout 的语义

1. init_process_group(timeout=) 约束什么

torch/distributed/distributed_c10d.py 的 init_process_group(timeout=None)：NCCL 后端默认 default_pg_nccl_timeout = 10 分钟（ProcessGroupNCCL.hpp 的 kProcessGroupNCCLDefaultTimeout = 10 × 60 × 1000 ms；其他后端 kProcessGroupDefaultTimeout = 30 分钟）。这个值成为每个 WorkNCCL 的 opTimeout_，也可以在 new_group(timeout=) 里按 PG 覆盖。

它约束的是：从 WorkNCCL 对象创建（即 CPU 线程调用 dist.all_reduce 的那一刻，workStartTime_ = steady_clock::now()）到 watchdog 观察到它的结束 event 完成之间的墙钟时间。把 CPU 线程、NCCL stream 与 watchdog 三条时间线并排，这个区间覆盖了什么、不覆盖什么就清楚了：

t0 = CPU 调用 dist.all_reduce：WorkNCCL 创建，workStartTime_ 起算
│
│ CPU 线程    ─┬─ 立即返回 ── Python 继续 ── .item()/wait() 阻塞 ─────────────▶
│              │                              （CPU 侧等多久，timeout 不管）
│ NCCL stream ─┴─[排在前面的 GEMM 30 s]─[NCCL kernel 自旋等对端 ……]─[end event]
│                                                                            │
│ watchdog     每 1 s：now − workStartTime_ > opTimeout_ 且 end event 未完成？
│
├──────────── opTimeout_ 计量区间（默认 10 min，含排队与等对端）──────────────┤
t0                                                              end event 完成

注意三点：

• 时钟从 CPU enqueue 开始，不是从 kernel 开始执行。如果 NCCL stream 前面排着一个 30 秒的 GEMM（第五篇的 stream 依赖），这 30 秒也算在里面；
• 它不约束 CPU 侧的任何等待。work.wait() 默认是 stream 级等待不阻塞 CPU（第五篇），Python 线程卡在别处与这个 timeout 无关——除非 TORCH_NCCL_BLOCKING_WAIT=1，那时 wait() 会在 CPU 上轮询 isCompleted() 并在同一个 timeout 后抛异常（这个模式下 2.12 不再创建 watchdog 线程，日志里会有 “TORCH_NCCL_BLOCKING_WAIT is enabled, NO watchdog thread is created”）；
• 它不是 NCCL 的机制。libnccl.so 里没有任何 timeout（IB 层的 NCCL_IB_TIMEOUT 是 QP 重传超时，只对物理丢包起作用），全部超时逻辑在 c10d 的 watchdog 线程里。用 PyNccl（vLLM）或裸 NCCL 的程序没有这层保护，hang 就是永远 hang——这是 RAS 存在的一个理由。

2. 为什么 checkpoint 保存或数据加载卡住会表现为 NCCL timeout

设想 rank 5 在第 3000 步之后保存 checkpoint，torch.save 因为共享存储抖动卡了 15 分钟。其他 63 个 rank 已经进入第 3001 步，调用了 all_reduce，它们的 WorkNCCL 从此刻计时。rank 5 还没调用第 3001 步的 all_reduce——它甚至还没开始第 3001 步的 forward。10 分钟后，63 个 rank 的 watchdog 同时判定它们的 all_reduce 超时，每个 rank 都打出 Watchdog caught collective operation timeout … OpType=ALLREDUCE，然后抛 DistBackendError（默认的 TORCH_NCCL_ASYNC_ERROR_HANDLING=3，即 SkipCleanUp 模式，不 abort communicator 直接退出；设为 1 的 TearDown 模式才会先 abort）。rank 5 十五分钟后从 torch.save 回来，发现 TCPStore 上有 dump 信号、或者自己调用 all_reduce 时 communicator 已被对端 abort，也报错退出。

事后看日志：64 个 rank 中 63 个报了 NCCL timeout，rank 5 报的是别的错（或者最后才报），所有报错都提到 all_reduce。根本原因是存储，现象是通信。这是 NCCL timeout 最常见的误导：timeout 只说明”有 rank 在这次集合通信上等了 10 分钟”，不说明为什么等——等的对象可能根本不在通信代码里。Flight Recorder 对这种情况给出的签名是上一章说的：rank 5 比别人少一条，且它的最后一条是 completed。

这也解释了为什么不能靠”调大 timeout”解决问题：timeout 是检测器，不是原因。把它调到 2 小时只会让 64 张卡多空转 2 小时。合理的做法反而是适当调小（比如训练稳态下一步不超过 30 秒，timeout 设 5 分钟），让问题更早暴露，同时用 Flight Recorder 保证暴露时有足够的信息。

3. heartbeat monitor：监视 watchdog 的人

watchdog 自己也可能卡——它调用 cudaEventQuery，在某些驱动错误下这个调用不返回；它在 timeout 处理里调用 ncclCommAbort，这个调用在网络异常时可能阻塞很久。HeartbeatMonitor（TORCH_NCCL_ENABLE_MONITORING，2.12 默认 true）每隔一段时间检查 watchdog 的心跳计数器，超过 TORCH_NCCL_HEARTBEAT_TIMEOUT_SEC（默认 480 秒）没变化就 dump（如果还能 dump）然后 std::abort() 杀掉本进程——宁可让作业死掉释放资源，也不让它无限期地占着 64 张卡。

时间线合起来（示意；dump 由 watchdog 广播信号、monitor 线程执行，heartbeat 只监视 watchdog 自身的活性）：

t = 0            rank 5 卡在 torch.save；其他 rank 调用 step 3001 的 all_reduce，WorkNCCL 计时开始
t = 10 min       63 个 rank 的 watchdog：checkTimeout 为真 → 打日志 → 广播 exception_dump → 等 60 s
t = 10 min + 1 s 各 rank 的 monitor 线程看到信号 → dumpDebuggingInfo → 写 trace_<rank>（rank 5 也写，它的 monitor 线程活着）
t = 11 min       63 个 rank 抛 DistBackendError；SkipCleanUp 模式下不 abort comm 直接退出
t = 15 min       rank 5 从 torch.save 回来，发现对端已退出，报错退出
若 watchdog 自己卡住：t = 10 min + 8 min，monitor 线程 abort 进程

4. 回答核心问题

一个 64 卡训练任务在第 3000 步 hang 住，所有 rank 的日志都停在 all_reduce。是谁的问题、是哪一次 all_reduce、为什么会等到 timeout 才暴露？

为什么等到 timeout 才暴露：因为 NCCL 的 kernel 在 flag 上自旋而没有超时概念，NCCL 库本身不区分”对端慢”与”对端不来”，CPU 线程在 enqueue 后早已返回。系统里唯一的计时器是 c10d watchdog 上每个 work 的 opTimeout_，默认 10 分钟。在这 10 分钟内没有任何组件有理由报错。”所有 rank 的日志都停在 all_reduce”只是因为 all_reduce 是每一步都要经过的同步点：不管谁掉队、因为什么掉队，其他人一定会停在下一次 all_reduce 上等他。

是谁的问题：分三步。第一步数进程：ps/Slurm 看 64 个进程是否都在，少了的那个就是答案（类 E-崩溃），去看它的 stderr 与 dmesg。第二步看 Flight Recorder：timeout 时每个 rank 已自动 dump 到 TORCH_NCCL_DEBUG_INFO_TEMP_FILE 前缀 + rank 号，跑 fr_trace.py。如果它报出某个序号上的 SIZE_OR_SYNTAX_MISMATCH / COLLECTIVE_TYPE_MISMATCH 并给出 culprit，那是代码问题（类 A/B），culprit 就是掉队者——去看它那条 entry 的 Python 栈，通常是一个只在特定 rank 或特定数据上走的分支。如果某个 rank 比别人少一条且它的最后一条是 completed，它卡在通信之外（类 E-卡住），对它 py-spy dump。如果 64 份记录完全一致，都停在同一序号且都是 started，那不是任何 rank 的代码问题，是网络或硬件（类 F）：查每个节点的 dmesg（Xid、mlx5 错误）、ibstat（State/Rate）、NCCL WARN 日志（Got completion … with error），交换机侧看端口计数。第三步——如果 Flight Recorder 没开（TORCH_NCCL_TRACE_BUFFER_SIZE=0，或用的是默认关闭的版本），退而用 py-spy dump 全体比栈，或者事后在复现时开 NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=COLL 比各 rank 最后一行的 opCount。

是哪一次 all_reduce：Flight Recorder 的 entry 有 collective_seq_id、profiling_name、input_sizes、pg_name 与调用时的 Python 栈。序号告诉你这是本 PG 的第几次集合通信；形状与 PG 名告诉你它是 DDP 的哪个 bucket、FSDP 的哪一层的 reduce_scatter、还是 TP 组的某次 all_reduce；Python 栈直接给出文件与行号。watchdog 日志里的 PG status: last enqueued work: N, last completed work: K 也给出同样的序号信息（FR dump 的 pg_status 里还有 last_started_collective），粒度粗一点。没有 Flight Recorder 时，NCCL_DEBUG_SUBSYS=COLL 的 opCount 是同一个东西的 NCCL 侧编号（每个 communicator 独立计数），配合 count 与 datatype 也能反推出是哪个 bucket。

“第 3000 步”这个信息本身也有用：如果 hang 稳定出现在同一步，几乎肯定是代码路径（第 3000 步做了 checkpoint / eval / 学习率调整 / 数据集切换）而不是网络；如果每次 hang 的步数不同，优先怀疑环境。

八、正确性问题

1. 浮点归约顺序：同一份数据、不同的和

浮点加法不满足结合律。ring all_reduce 里第 k 个分块在 rank k 上按 ring 的顺序累加，tree 按树的结构累加，NVLS 在 NVSwitch 里按硬件的顺序累加；同一组输入，三种算法给出的 sum 在最后几位上可以不同。同一算法内部，channel 数、chunk 大小的变化也会改变累加顺序。所以：

• 同一作业内两次运行结果的 bit 级差异，如果只发生在最后几个有效位、且随算法/规模变化，是正常的；
• 各 rank 拿到的结果互相一致是 all_reduce 的语义保证（每个分块只在一处归约再广播），不一致是 bug（见第 4 小节的数据竞争）；
• 需要 bit 级可复现时，固定 NCCL_ALGO、NCCL_PROTO、channel 数，并接受性能损失；或者在框架层用确定性归约（例如先 all_gather 再本地按固定顺序求和，通信量增加 n 倍）。

nccl-tests -c 1 的校验绕开了这个问题：src/common.cu 的 CheckDelta 调用 verifiable/verifiable.cu 的 ncclVerifiableVerify，输入数据由 ncclVerifiablePrepareInput 按 rank 数、rank 号与元素下标精心构造（verifiable.h 的注释：carefully craft floating point input to produce exactly predictable output），使归约结果不依赖累加顺序，所以浮点类型也做精确比对，只有 avg 算子对浮点类型放宽到 2 ulp（因为预乘 1/n 的中间精度未定义）。因此 #wrong 非零基本可以视为真正的数据错误，而不是归约顺序引起的精度差异。

2. 算法依赖的数值差异

LL 和 LL128 协议在传输时把数据与 flag 打包（LL 8 字节数据 + 8 字节 flag，LL128 每 128 字节 120 字节数据），Simple 协议整块传输，三者的数值结果应当相同——协议只改变搬运方式，不改变累加顺序。算法（Ring / Tree / NVLS / PAT）会改变累加顺序。所以调试”换了 NCCL 版本 loss 曲线不一样了”时，先用 NCCL_ALGO 固定算法看差异是否消失；消失则是归约顺序，是正常的；不消失再怀疑 bug。

历史上确有 LL128 在特定 PCIe 拓扑上因为 128 字节原子性假设不成立而产生错误数据的案例，NCCL 在拓扑探测里已针对已知平台禁用；如果 -c 1 在某台机器上报 #wrong，NCCL_PROTO=^LL128 对照一次是标准动作。

3. NaN 的来源与 TORCH_NCCL_NAN_CHECK

一个 rank 的梯度里出现 NaN，all_reduce 之后所有 rank 的梯度都是 NaN——归约把 NaN 传染给所有人，事后无法知道源头是谁。TORCH_NCCL_NAN_CHECK=1（ProcessGroupNCCL.cpp 的 enableNanCheck_）让 ProcessGroupNCCL 在每次集合通信与 send 之前在 NCCL stream 上跑一个 checkForNan kernel（torch/csrc/distributed/c10d/NanCheck.cu）检查输入；发现 NaN 就在发送方报错，错误里带 rank 与操作名。它有开销（每次通信多一个扫描 kernel，约等于多读一遍输入），只在追 NaN 时开。

它检查的是输入不是输出，所以能区分”我这边产生了 NaN”与”别人传给我 NaN”：前者本 rank 报错，后者本 rank 不报、别的 rank 报。NaN 的来源通常在计算侧（溢出、除零、坏数据），通信层只是传播者；但有一类通信层引入的 NaN 值得知道——第 4 小节的数据竞争让 all_reduce 读到了尚未写完的 buffer，读到的垃圾恰好是 NaN 模式。

4. 多 communicator 与多 stream 的数据竞争

第五篇讲了 ProcessGroupNCCL 用自己的 stream，与计算 stream 靠 event 建依赖，输入 tensor 的生命周期靠 recordStream 或 stash 保护。当用户绕过这套机制时会出问题：

• 自建 stream 上调用集合通信而不建依赖：with torch.cuda.stream(s): dist.all_reduce(t)——ProcessGroupNCCL 会让 NCCL stream 等待当前 stream s，但如果 t 是在默认 stream 上刚算出来的、而 s 与默认 stream 之间没有 event，all_reduce 可能读到旧值。症状是结果偶尔错、与负载相关。
• 两个 PG 在两条 NCCL stream 上并发操作重叠的 buffer：DP 组 all_reduce 的梯度 buffer 与 TP 组 all_reduce 的激活 buffer 不会重叠，但用户自己 new_group 出的多个组如果操作同一个 tensor 的不同 view 且不同步，会竞争。
• async_op=True 之后在 wait() 之前修改输入：第五篇核心问题的答案——修改会被通信 kernel 读到，结果不确定。
• 输入 tensor 被 Caching Allocator 提前回收再分配给别的计算：TORCH_NCCL_AVOID_RECORD_STREAMS 的 stash 机制（2.12 已是默认，再设该变量只会得到一条 deprecated 警告）防的就是这个；自写扩展直接调 NCCL 时没有这层保护。

这类问题的共同签名是：各 rank 的 all_reduce 结果不一致（用 all_gather 收齐来比，或者 TORCH_DISTRIBUTED_DEBUG=DETAIL 让 c10d 对每次集合通信的输入做哈希并跨 rank 比对——enableCollectiveHashDebug_），并且随时序变化。定位手段是 CUDA_LAUNCH_BLOCKING=1 或强制在通信前后 torch.cuda.synchronize()——问题消失即证实是竞争，然后二分找缺失的依赖。

5. 用 nccl-tests 做正确性验收

新机器、新驱动、新 NCCL 版本上线前跑一遍带校验的全量扫描：

for t in all_reduce reduce_scatter all_gather broadcast reduce alltoall sendrecv; do
  ./build/${t}_perf -b 8 -e 1G -f 2 -g 8 -n 5 -w 2 -c 1 -o all -d all -T 600 2>&1 | tee check.${t}.log
done
grep -H "Out of bounds" check.*.log     # 全部应为 "0 OK"

-o all -d all 让每个算子与每种类型的组合都跑一遍。它不能发现框架层的竞争，但能排除 NCCL 与硬件层的数据错误（坏的 NVLink 链路、PCIe 的 relaxed ordering 问题、GDR 缓存一致性问题都曾以 #wrong 非零的形式被发现）。

九、决策树

把前面各章收成三棵树。每个叶子给出”看什么”与”做什么”。

1. 慢

慢
├─ 先做：nccl-tests 单机基线 + 多机基线 + 与第三章参考线比 + 与 cost_model 预测比
│
├─ 大消息 busbw 低（节点内）
│    ├─ 查 nvidia-smi topo -m：GPU 间是 NV# 还是 PIX/PXB/SYS？NVLink 是否降级（nvidia-smi nvlink -s）
│    ├─ 查 NCCL_DEBUG=INFO："via P2P/…" 还是 "via SHM/…"？  → SHM 说明 P2P 被禁或 level 不够：查 NCCL_P2P_LEVEL/DISABLE、IOMMU、ACS
│    ├─ 查 channel 数（TUNING 日志的 channel{Lo..Hi} 或 "Channel NN/…" 最大编号） → 明显少于预期：试 NCCL_MAX_NCHANNELS 对照，怀疑拓扑探测（NCCL_TOPO_DUMP_FILE 看 XML）
│    └─ H100 上低于 350 GB/s 且无 "NVLS multicast support is available"  → 驱动/Fabric Manager
│
├─ 大消息 busbw 低（跨机）
│    ├─ 先单独测网络：ib_write_bw 两机之间能到多少？低 → 网络问题，不是 NCCL 的
│    ├─ "Using network Socket"  → 没走 RDMA：libibverbs/驱动/NCCL_IB_DISABLE/容器里缺 /dev/infiniband
│    ├─ "GPU Direct RDMA Disabled … distance N > M"  → GDR 没开：nvidia-peermem/DMA-BUF 是否加载；GPU 与网卡是否同 PCIe switch；修亲和而不是放宽 level
│    ├─ "NET/IB : Using [...]" 列出的设备数 < GPU 数，或混入了不同链路类型  → NCCL_IB_HCA 限定
│    ├─ "via NET/IB/<dev>" 的 dev 与 GPU 不亲和  → 拓扑探测或容器拓扑可见性问题；NCCL_TOPO_FILE 兜底
│    ├─ RoCE：ibstat 的 port_rcv_errors / 交换机丢包计数上升  → PFC/ECN、NCCL_IB_TC 与交换机队列不一致
│    ├─ 带宽不稳、时高时低  → ECMP 哈希不均：NCCL_IB_QPS_PER_CONNECTION=2/4；或 rail 网络上 NCCL_CROSS_NIC=0
│    └─ 都正常但仍差 20%+  → NCCL_BUFFSIZE（高延迟链路）；NCCL 版本；固件
│
├─ 小消息延迟高（time 列左端偏大）
│    ├─ NCCL_DEBUG_SUBSYS=TUNING 看 proto：小消息应为 LL/LL128；是 Simple → 查 NCCL_PROTO 是否被设置、LL128 是否被拓扑禁用
│    ├─ 进程绑核：numactl / taskset 是否让进程与 GPU、网卡在同一 NUMA
│    ├─ proxy 线程被抢占：top 看 CPU 是否满载；训练时预留核给 NCCL
│    └─ 训练里小消息多而 nccl-tests 正常  → 框架没合并：DDP bucket、FSDP 的 reshard 策略、functional collective 的重排
│
└─ 多机远差于单机 / 节点数
     └─ 跨机的每一项 + 检查是否每节点网卡数 < GPU 数 + NCCL_CROSS_NIC

2. hang

hang（watchdog 报 timeout，或没有 watchdog 时进程不动）
├─ 先做：数进程数；收集全体 py-spy dump；收集 Flight Recorder dump；grep 全体 NCCL WARN；看 dmesg
│
├─ 进程数少了  → 类 E-崩溃：看缺失 rank 的 stderr（OOM/异常）、dmesg（Killed process / Xid）
│               处理：修根因；开 TORCH_NCCL_PROPAGATE_ERROR 让其他 rank 早退
│
├─ 所有 rank 的 FR 记录一致（同序号、同状态 started）
│    ├─ 有 NCCL WARN（completion error / Cuda failure / socket closed）  → 类 F：网络或 GPU 硬件
│    │      ibstat（State/Rate）· dmesg（mlx5、Xid）· nvidia-smi -q（Xid、ECC）· 交换机
│    └─ 无 WARN  → 也可能是类 F 的静默丢失，或 NCCL bug；cuda-gdb 确认 kernel 在自旋；换 NCCL 版本对照
│
├─ FR 报某序号 SIZE_OR_SYNTAX / TYPE / DTYPE MISMATCH，给出 culprit
│    → 类 A/B：代码问题。看 culprit 那条 entry 的 Python 栈，找只在该 rank 上走的分支
│      （按 rank 的 if、按数据的 if、动态 shape、只在 rank 0 做的 eval/log）
│
├─ 某 rank 比别人少一条，且其最后一条是 completed
│    → 类 E-卡住：对它 py-spy dump。常见：torch.save、DataLoader、logging、GIL 死锁、第三方锁
│
├─ 部分 rank 停在 send/recv、部分停在集合通信；FR 里 is_p2p 条目找不到配对
│    → 类 C：流水线 stage 边界的 send/recv 数量或 peer 不一致；检查 micro-batch 数、warmup 阶段的调度
│
└─ 各 rank 停在不同 PG 的调用上
     → 类 D：多 communicator 顺序不一致。统一各 rank 上 PG 的调用顺序；或用 batch_isend_irecv / coalescing 合并

3. 错

错
├─ 各 rank 结果互相一致，但与预期 / 上一版本有微小差异
│    → 浮点归约顺序：固定 NCCL_ALGO 看差异是否消失；正常现象，需要复现性时固定算法与 channel
│
├─ 各 rank 结果互相不一致
│    ├─ TORCH_DISTRIBUTED_DEBUG=DETAIL 的输入哈希是否一致？不一致 → 输入就不同：数据竞争或调用不一致
│    ├─ 加 torch.cuda.synchronize() 后消失  → stream 依赖缺失 / async_op 后过早修改输入 / allocator 回收
│    └─ nccl-tests -c 1 也报 #wrong  → NCCL 或硬件层：NCCL_PROTO=^LL128 对照；nvidia-smi -q 看 ECC/Xid；换卡
│
└─ NaN
     ├─ TORCH_NCCL_NAN_CHECK=1 → 报错的 rank 就是源头；看该 rank 的计算（溢出、除零、坏样本）
     └─ 无 rank 报错但结果 NaN  → 通信过程中引入：数据竞争（见上）或硬件错误

4. 一小时排障流程

第 0–10 分钟   收集：进程数、全体 py-spy dump、FR dump、NCCL WARN、dmesg、ibstat、nvidia-smi -q
第 10–20 分钟  分类：按上面三棵树的第一层判断是 慢/卡/错，是 代码/环境
第 20–40 分钟  定位：fr_trace.py 对齐（卡）/ nccl-tests 基线与 INFO 日志八行（慢）/ 固定算法对照（错）
第 40–60 分钟  验证：改一处、跑 nccl-tests 或 hang_lab 复现脚本确认；把结论与环境记录归档

超过一小时还没有定位到层，通常是因为跳过了第 0–10 分钟的收集直接开始猜。

十、本文小结

1. 要点回顾

nccl-tests      总 rank 数 = 进程 × -t × -g；-b/-e/-f 扫描，-n/-w 迭代预热（-w 默认仅 1），-c 0 关校验测性能；algo/proto/channel 看 NCCL 的 TUNING 日志
                busbw = algbw × 系数（all_reduce 2(n-1)/n，all_gather/reduce_scatter/all_to_all (n-1)/n，broadcast 1），
                以各 src/*.cu 的 *GetBw 为准；只有 busbw 能和硬件比
曲线            左端斜线高度看 α（直接看 time 列），右端平台看 β，拐点 S_knee = n·α·β，到 90% 平台约 9·S_knee
                8×H100 节点内 ≈ 10 MB 拐点、~90 MB 接近平台；64 卡跨机 ≈ 30 MB 拐点、~260 MB 接近平台（典型值，非实测）
参考线          8×H100 NVSwitch 350–480 GB/s（NVLS 可超单向 450）；8×A100 220–280；PCIe-only 10–25；
                2 节点 NDR 每 GPU 40–48；HDR 20–24；每节点 1 张 NDR 5–6（全部非实测）
参数            NCCL_PARAM 整数、只读一次、优先级 env > NCCL_CONF_FILE > ~/.nccl.conf > /etc/nccl.conf > 默认；
                字符串类走 ncclGetEnv；P2P_LEVEL/NET_GDR_LEVEL 走 ncclGetLevel 接受路径类型名
                应设：IB_HCA、SOCKET_IFNAME、GID_INDEX/TC（RoCE）、DEBUG、DEBUG_FILE；
                对照用：ALGO/PROTO、P2P/SHM/GDR、NCHANNELS、CROSS_NIC、QPS；几乎不动：NTHREADS、BUFFSIZE、IB_TIMEOUT
日志            WARN 生产、INFO 排障（默认子系统 INIT,BOOTSTRAP,ENV）、TRACE 需 -DENABLE_TRACE；
                SUBSYS 选 GRAPH/NET/TUNING/COLL；DEBUG_FILE 的 %h/%p；八行速查；RAS/ncclras 是 NCCL 侧的 opCount 对齐
hang            kernel 自旋无超时、CPU 早已返回、唯一计时器是 c10d watchdog 的 opTimeout_（默认 10 min，从 enqueue 起算）
                六类：参数不一致 / 次数不一致 / send-recv 不配对 / 多 comm 交叉 / rank 崩或卡在别处 / 网络硬件；
                前四类各 rank 最后一次操作不一致，后两类一致
工具            py-spy dump 看 Python 栈 → gdb 看 C++ → cuda-gdb 看 kernel；
                Flight Recorder 记每次调用（seq、形状、dtype、栈、状态），timeout 时经 TCPStore 通知全体 dump，
                fr_trace.py 按 collective_seq_id 对齐给出 MatchState 与 culprit；DESYNC_DEBUG 是轻量前身
timeout         约束 enqueue → 结束 event 的墙钟时间；不约束 CPU；不是 NCCL 的机制；
                checkpoint/dataloader 卡住 → 其他 rank 在下一次集合通信上等满 timeout → 表现为 NCCL timeout
正确性          归约顺序随算法/channel 变，各 rank 间一致是语义保证；NAN_CHECK 在发送方检查输入；
                结果不一致 + 随时序变 = stream 竞争；nccl-tests -c 1 -o all -d all 做硬件层验收

2. 排障检查项

通信层出问题时按这个顺序看：

3. 本篇涉及的源码与工具位置

项目          路径                                                        关键函数 / 内容
─────────────────────────────────────────────────────────────────────────────────────────────────────────────
nccl-tests    src/common.cu                                               getopt_long 参数表与 usage；参数默认值；run()；MPI 的使用；NCCL_TESTS_MIN_BW；
2.18.3                                                                    NCCL_TESTS_DEVICE；CheckDelta
              src/common.h                                                NCCL_TESTS_VERSION
              src/util.cu                                                 writeResultHeader 列名；writeResultFooter 的 "Out of bounds values" / "Avg bus bandwidth"
              src/all_reduce.cu 等                                        *GetBw：algbw → busbw 的系数
              verifiable/verifiable.cu                                    ncclVerifiablePrepareInput / ncclVerifiableVerify（精确校验；avg 放宽 2 ulp）
              README.md · doc/PERFORMANCE.md                              编译选项 MPI=1；Arguments 与默认值；busbw 的推导
NCCL 2.28.9   src/include/param.h · src/misc/param.cc                     NCCL_PARAM 宏；ncclLoadParam；initEnv 与 NCCL_CONF_FILE；ncclGetEnv
              src/plugin/env.cc                                           NCCL_ENV_PLUGIN
              src/debug.cc · src/include/debug.h                          ncclDebugInit：NCCL_DEBUG 级别、NCCL_DEBUG_SUBSYS、NCCL_DEBUG_FILE 的 %h/%p；
                                                                          ncclDebugLog 行格式；TRACE 需 ENABLE_TRACE
              src/graph/tuning.cc                                         NCCL_ALGO / NCCL_PROTO 的 parseList；NCCL_NTHREADS；TUNING 日志
              src/graph/connect.cc                                        NCCL_MIN_NCHANNELS / NCCL_MAX_NCHANNELS
              src/graph/paths.cc                                          ncclGetLevel；NCCL_P2P_LEVEL / NCCL_P2P_DISABLE / NCCL_NET_GDR_LEVEL；
                                                                          "GPU Direct RDMA Enabled/Disabled"
              src/graph/topo.cc                                           topoPathTypeStr（LOC NVL … SYS）；NCCL_TOPO_FILE
              src/graph/search.cc                                         NCCL_CROSS_NIC；NCCL_GRAPH_DUMP_FILE
              src/init.cc                                                 NCCL_BUFFSIZE；NCCL_COLLNET_ENABLE；NCCL_TOPO_DUMP_FILE；NCCL_NET；ncclAlgoStr/ncclProtoStr
              src/transport/net_ib.cc                                     NCCL_IB_HCA 解析（^ 与 =）；IB_GID_INDEX / IB_TC / IB_SL / IB_QPS_PER_CONNECTION /
                                                                          IB_SPLIT_DATA_ON_QPS / IB_TIMEOUT / IB_RETRY_CNT；"NET/IB : Using …"
              src/transport/net.cc · p2p.cc · shm.cc                     "via NET/IB/…/GDRDMA" · "via P2P/…" · "via SHM/…"
              src/transport/nvls.cc                                       NCCL_NVLS_ENABLE
              src/transport/net_socket.cc · src/misc/socket.cc            NCCL_SOCKET_NTHREADS / NCCL_NSOCKS_PERTHREAD；NCCL_SOCKET_IFNAME
              src/enqueue.cc                                              COLL 日志的 opCount 行；TUNING 日志的 "N Bytes -> Algo … proto …"
              src/ras/ · src/bootstrap.cc                                 RAS 子系统；ncclras 客户端（client.cc）；collOpCounts 比对（client_support.cc）；NCCL_RAS_ENABLE
PyTorch       torch/csrc/distributed/c10d/ProcessGroupNCCL.hpp            TORCH_NCCL_* 变量定义；ErrorHandlingMode；kStoreDumpKey；DumpPipe；
2.12                                                                      kProcessGroupNCCLDefaultTimeout
              torch/csrc/distributed/c10d/ProcessGroupNCCL.cpp            各 TORCH_NCCL_* 的 getCvar 默认值；WorkNCCL::checkTimeout；Watchdog::runLoop；
                                                                          HeartbeatMonitor::runLoop；broadcastDumpSignal；dumpDebuggingInfo；DesyncDebugger；
                                                                          checkForNan 调用点
              torch/csrc/distributed/c10d/FlightRecorder.hpp/.cpp         Entry 字段；DebugInfoWriter 默认路径；TORCH_FR_* 别名
              torch/csrc/distributed/c10d/FlightRecorderCuda.cpp          dump_nccl_trace
              torch/csrc/distributed/c10d/TraceUtils.h                    retrieveDesyncReport
              torch/csrc/distributed/c10d/NanCheck.hpp/.cu                checkForNan
              torch/distributed/flight_recorder/fr_trace.py · components/ 分析工具；MatchState；参数
              torch/distributed/distributed_c10d.py · constants.py        init_process_group(timeout=)；default_pg_nccl_timeout

4. comm-probe 本篇增量：sweep.sh 与 hang_lab/

sweep.sh 对一组原语跑 nccl-tests 扫描、解析成 CSV、画出 busbw 曲线并叠上第一篇 cost_model.py 的 ring 预测：

#!/usr/bin/env bash
# comm-probe/sweep.sh — nccl-tests 扫描 → CSV → 图
# 用法: sweep.sh [-c all_reduce,all_gather] [-g 8] [-b 8] [-e 8G] [-m "mpirun -np 16 -N 8 --hostfile hosts"]
#       环境: NCCL_TESTS_BIN=nccl-tests/build  ALPHA_US=3 BETA_GBS=400（预测线用）
set -euo pipefail
COLLS="all_reduce"; NGPUS=8; MINB=8; MAXB=8G; LAUNCHER=""; ITERS=20; WARM=5
while getopts "c:g:b:e:m:n:w:" o; do case $o in
  c) COLLS=$OPTARG;; g) NGPUS=$OPTARG;; b) MINB=$OPTARG;; e) MAXB=$OPTARG;;
  m) LAUNCHER=$OPTARG;; n) ITERS=$OPTARG;; w) WARM=$OPTARG;; esac; done
BIN=${NCCL_TESTS_BIN:-./nccl-tests/build}
OUT=${OUT:-sweep_$(hostname -s)_$(date +%Y%m%d_%H%M%S)}; mkdir -p "$OUT"
# 记录环境：没有这些，三个月后无法解释两条曲线为什么不同
{ nvidia-smi topo -m; nvidia-smi --query-gpu=name,driver_version --format=csv; env | grep '^NCCL_' || true; \
  ldd "$BIN/all_reduce_perf" | grep nccl; } > "$OUT/env.txt" 2>&1
[ -n "$LAUNCHER" ] && G=1 || G=$NGPUS          # mpirun 下每进程一卡
echo "coll,nranks,size_bytes,count,type,time_us,algbw_gbs,busbw_gbs,wrong" > "$OUT/sweep.csv"
for coll in ${COLLS//,/ }; do
  log="$OUT/$coll.log"
  NCCL_DEBUG=${NCCL_DEBUG:-WARN} $LAUNCHER "$BIN/${coll}_perf" -b "$MINB" -e "$MAXB" -f 2 -g "$G" \
      -n "$ITERS" -w "$WARM" -c 0 -T 600 2>&1 | tee "$log"
  # 数据行以数字开头；取 out-of-place 的 size count type time algbw busbw #wrong（第 1,2,3,6,7,8,9 列）
  nranks=$(grep -c '^#  Rank ' "$log" || echo "$NGPUS")
  awk -v c="$coll" -v n="$nranks" '$1 ~ /^[0-9]+$/ {printf "%s,%s,%s,%s,%s,%s,%s,%s,%s\n", c,n,$1,$2,$3,$6,$7,$8,$9}' \
      "$log" >> "$OUT/sweep.csv"
  grep -E "Out of bounds|Avg bus bandwidth" "$log" >> "$OUT/summary.txt" || true
done
python3 "$(dirname "$0")/plot_sweep.py" "$OUT/sweep.csv" --alpha-us "${ALPHA_US:-3}" --beta-gbs "${BETA_GBS:-400}" -o "$OUT/busbw.png"
echo "done: $OUT"

plot_sweep.py 读 CSV，按原语画 busbw 对 size 的对数曲线，并叠上 ring 预测（与 cost_model.py 的 ring_allreduce_time 公式相同，这里内联以便脚本独立运行）：

#!/usr/bin/env python3
# comm-probe/plot_sweep.py
import argparse, csv, collections
import matplotlib; matplotlib.use("Agg")
import matplotlib.pyplot as plt

def ring_allreduce_busbw(n, S, alpha, beta):
    """第 1 篇 cost_model.py 的 ring 公式：T = 2(n-1)α + 2(n-1)/n · S/β；busbw = 2(n-1)/n · S / T"""
    T = 2 * (n - 1) * alpha + 2 * (n - 1) / n * S / beta
    return 2 * (n - 1) / n * S / T

ap = argparse.ArgumentParser()
ap.add_argument("csv"); ap.add_argument("-o", default="busbw.png")
ap.add_argument("--alpha-us", type=float, default=3.0); ap.add_argument("--beta-gbs", type=float, default=400.0)
a = ap.parse_args()
rows = collections.defaultdict(list)
with open(a.csv) as f:
    for r in csv.DictReader(f):
        rows[(r["coll"], int(r["nranks"]))].append((int(r["size_bytes"]), float(r["busbw_gbs"]), float(r["time_us"])))
fig, ax = plt.subplots(figsize=(9, 5))
for (coll, n), pts in rows.items():
    pts.sort(); xs = [p[0] for p in pts]
    ax.plot(xs, [p[1] for p in pts], marker="o", ms=3, label=f"{coll} n={n} (measured)")
    if coll == "all_reduce":
        pred = [ring_allreduce_busbw(n, S, a.alpha_us * 1e-6, a.beta_gbs * 1e9) / 1e9 for S in xs]
        ax.plot(xs, pred, ls="--", label=f"ring model α={a.alpha_us}µs β={a.beta_gbs}GB/s")
        knee = n * a.alpha_us * 1e-6 * a.beta_gbs * 1e9
        ax.axvline(knee, color="gray", ls=":", lw=1); ax.text(knee, 1, f"S_knee≈{knee/2**20:.0f}MiB", rotation=90, va="bottom")
ax.set_xscale("log", base=2); ax.set_xlabel("message size (bytes)"); ax.set_ylabel("busbw (GB/s)")
ax.grid(True, which="both", alpha=.3); ax.legend(); fig.tight_layout(); fig.savefig(a.o, dpi=120)

预测线与实测线的关系按第三章第 2 小节读：平台差是 β，左端差是 α，中段差是切换点；要对照台阶位置，把 NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=TUNING 传给 sweep.sh（它只在 NCCL_DEBUG 未设时默认 WARN），日志里 rank 0 的 N Bytes -> Algo … proto … 行按字节数与结果表逐行对应。

hang_lab/ 是三个 torchrun 脚本加一份诊断剧本。三个脚本共用一个骨架，每步的 tensor 大小随步数变化（让 Flight Recorder 的匹配无歧义），并把 timeout 设短：

# comm-probe/hang_lab/common.py
import os, datetime, torch, torch.distributed as dist

def setup(timeout_s=60):
    dist.init_process_group("nccl", timeout=datetime.timedelta(seconds=timeout_s))
    rank = dist.get_rank(); torch.cuda.set_device(rank % torch.cuda.device_count())
    return rank, dist.get_world_size()

def numel(step):            # 每步大小不同：(step+1) × 64K 元素
    return (step + 1) << 16

# comm-probe/hang_lab/size_mismatch.py — 类 A：rank 3 在第 5 步多传 1024 个元素
from common import setup, numel
import torch, torch.distributed as dist
rank, world = setup()
for step in range(8):
    n = numel(step) + (1024 if (step == 5 and rank == 3) else 0)
    t = torch.ones(n, device="cuda")
    dist.all_reduce(t)                 # 第 5 步：rank 3 的 count 与其他 rank 不一致 → 通常 hang
    torch.cuda.synchronize()
    if rank == 0: print(f"step {step} done", flush=True)
dist.destroy_process_group()

另两个脚本只改一行：count_mismatch.py 在第 5 步让 rank 2 跳过 all_reduce（if not (step == 5 and rank == 2): dist.all_reduce(t)），之后它的第 6 次调用与别人的第 5 次配对，形状不同；unpaired_sendrecv.py 在每步 all_reduce 之后加一段流水线式的 send/recv（rank i 发给 i+1），并在第 5 步让 rank 1 不 recv，于是 rank 0 的 send 永远等不到配对，rank 0 与 rank 1 先停，其余 rank 在第 6 步的 all_reduce 上被连锁拖住。

统一的启动与诊断剧本：

# 启动：开 Flight Recorder，dump 到当前目录，开命名管道以便不等 timeout 就 dump
export TORCH_NCCL_TRACE_BUFFER_SIZE=2000 TORCH_NCCL_DUMP_ON_TIMEOUT=1 TORCH_NCCL_ENABLE_MONITORING=1
export TORCH_NCCL_DEBUG_INFO_TEMP_FILE=$PWD/fr/trace_  TORCH_NCCL_DEBUG_INFO_PIPE_FILE=$PWD/fr/pipe
export NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=COLL NCCL_DEBUG_FILE=$PWD/fr/nccl.%h.%p.log
mkdir -p fr; torchrun --nproc_per_node 8 hang_lab/size_mismatch.py &

# 观察到 "step 4 done" 之后不再有输出（约 5 秒后），开始诊断：
# 1. 进程都在吗？                                → 8 个都在，排除崩溃
pgrep -f hang_lab | wc -l
# 2. 全体 Python 栈                              → 都停在 all_reduce 那一行；栈相同 → 不是"卡在别处"
for p in $(pgrep -f hang_lab/size_mismatch); do echo "== $p"; py-spy dump --pid $p | head -20; done
# 3. NCCL 侧 opCount 对齐                        → 各 rank 最后一行 opCount 相同（同一次调用），count 列 rank 3 多 1024
for f in fr/nccl.*.log; do echo "$f: $(grep 'AllReduce: opCount' $f | tail -1 | grep -o 'opCount [0-9a-f]* .* count [0-9]*')"; done
# 4. 不等 60 s timeout，手动触发 FR dump（DumpPipe 的文件名 = 前缀 + 全局 rank + ".pipe"）
for r in 0 1 2 3 4 5 6 7; do echo 1 > fr/pipe${r}.pipe; done; sleep 3; ls fr/trace_*
# 5. 对齐
python -m torch.distributed.flight_recorder.fr_trace fr/ -p trace_

三种情形下 fr_trace.py 的预期输出与判读：

size_mismatch.py       collective_seq_id=6（第 5 步，从 1 起算）：SIZE_OR_SYNTAX_MISMATCH，culprit rank 3；
                       该条 input_sizes 为 [394240] 而其他为 [393216]；--print_stack_trace 指到 dist.all_reduce 那一行
                       → 类 A。修法：在 all_reduce 前 assert 形状（或在开发期开 TORCH_DISTRIBUTED_DEBUG=DETAIL 做跨 rank 形状校验）

count_mismatch.py      collective_seq_id=6：rank 2 的 input_sizes 是 [458752]（第 6 步的大小），其他是 [393216]
                       → SIZE_OR_SYNTAX_MISMATCH，culprit rank 2；序号 6 上它"看起来"传了别的东西，实际是少调了一次
                       py-spy 会看到 rank 2 停在 step 6 的 all_reduce、其他 rank 停在 step 5
                       → 类 B。修法：找出让 rank 2 跳过的分支；集合通信必须在所有 rank 上无条件执行
                       （对照实验：把 numel() 改成常数，fr_trace 会一路 FULLY_MATCHED 直到最后一个序号才报 STATE_MISMATCH——
                       这是"每步大小相同时定位不到具体哪一步"的直观演示）

unpaired_sendrecv.py   集合通信条目到 seq 5 全部 FULLY_MATCHED；p2p 条目里 rank 0 有一条 send(dst=1) 无配对，rank 1 无对应 recv；
                       rank 2–7 停在 seq 6 的 all_reduce（等 rank 0、1），rank 0 停在 send，rank 1 停在下一步的 all_reduce
                       py-spy 直接显示两组不同的栈（send vs all_reduce）
                       → 类 C。修法：send/recv 成对写在同一段代码里（batch_isend_irecv），stage 边界的 micro-batch 数各 rank 一致

每个脚本都在 60 秒 timeout 后自动 dump 并退出，剧本第 4 步只是为了不等。跑完三个脚本，读者应当对”各 rank 最后一次操作一致还是不一致”这个判据、以及 py-spy 与 Flight Recorder 各自能回答什么问题有肌肉记忆——这是第九章决策树 hang 分支的全部基础。

到这里，训练侧的通信从理论到测量到排障已经闭环。下一篇转向推理：decode 阶段每层一次几十 KB 的 TP all_reduce，纯延迟主导，NCCL 的固定开销成了主要成本；PD 分离的 KV 传输是点对点大块搬运，不需要归约也不需要 NCCL。两个场景都要用本篇的方法测、用第一篇的模型算，但答案会不一样。

8 卡 TP 的 decode，每层一次 128 KB 的 all_reduce，NCCL 要 30 微秒，custom all-reduce 要 10 微秒。这 20 微秒省在哪里？为什么这个方法不能用在训练的梯度同步上？

十一、自测

1. mpirun -np 8 all_reduce_perf -b 8 -e 8G -f 2 -g 1 -w 5 -n 20 -c 0 各参数什么意思？总 rank 数是多少？

   答案

   8 个进程、每进程 1 张 GPU（-g 1）→ 8 个 rank；从 8 B 到 8 GB 每次 ×2 扫描；预热 5 次、测 20 次；-c 0 关掉结果校验只测性能。总 rank = 进程 × -t × -g。

2. nccl-tests 的曲线左端斜线、右端平台、拐点各由什么决定？8×H100 节点内的拐点与接近平台大约在哪？

   答案

   左端由 \(\alpha\) 决定（直接看 time 列）、右端平台由 \(\beta\) 决定、拐点 \(S_{knee} = n\alpha\beta\)、到 90% 平台约 \(9 S_{knee}\)；8×H100 节点内约 10 MB 拐点、约 90 MB 接近平台（典型值）。

3. 8×H100 NVSwitch 机器 all_reduce 的 busbw 平台该在什么范围？测出 150 GB/s 该怀疑什么？

   答案

   350–480 GB/s（NVLS 可超单向 450）；150 太低——查是否落到了 SHM / PCIe 路径（NCCL_DEBUG=INFO 看 transport）、NVLink 有无 inactive 链路、channel 数被限、或 GPU 时钟 / 功耗被限。

4. NCCL 参数的优先级顺序是什么？哪几个参数是“应设”、哪几个“几乎不动”？

   答案

   env > NCCL_CONF_FILE > ~/.nccl.conf > /etc/nccl.conf > 默认，整数参数只读一次。应设：IB_HCA、SOCKET_IFNAME、RoCE 的 GID_INDEX / TC、DEBUG / DEBUG_FILE；对照用：ALGO / PROTO、P2P / SHM / GDR、NCHANNELS；几乎不动：NTHREADS、BUFFSIZE、IB_TIMEOUT。

5. 一个 rank 在写 checkpoint 时卡了 15 分钟，最终日志报的是什么错？为什么？

   答案

   其他 rank 报 NCCL timeout（opTimeout_ 10 分钟）——它们在下一次集合通信上等这个 rank，watchdog 到期直接退出；卡住的 rank 本身可能没有任何 NCCL 日志。timeout 约束的是 enqueue 到 end event 的墙钟，不区分原因，所以 NCCL timeout 常是别处 hang 的表现。

下一篇

推理侧的通信：custom all-reduce 与 KV 传输

• 上一篇
  通信与互联（05）：PyTorch 的通信栈——ProcessGroupNCCL、stream 语义与计算通信重叠
• 下一篇
  通信与互联（07）：推理侧的通信——custom all-reduce 与 KV 传输