本章问题 训练在第一次 collective 之前 hang 住时，“NCCL 初始化失败”几乎没有诊断价值。 初始化至少包含： 参数和 CUDA runtime 校验 unique ID / bootstrap root 全部 rank rendezvous bootstrap ring 建立 peerInfo AllGather local rank / node 划分 topolo...

本章问题 不设置 NCCL_ALGO 和 NCCL_PROTO 时，NCCL 如何从 Tree、Ring、LL、 LL128 和 Simple 中选出一个组合？ “NCCL 会 benchmark 后选择最快实现”是错误答案。NCCL 2.22.3 在 communicator 初始化时根据 topology graph、硬件类型和经验常量建立 latency/bandwidth tab...

本章问题 NCCL 性能排障中最容易混用的词是： channel CUDA block / CTA worker thread / blockDim.x chunk slice step buffer loop 如果把它们都理解成“并行度”，就无法解释： 为什么请求 12 channels 的 4 KiB AllReduce 实际只 launch 1 个 CTA。 为什么...

本章问题 NCCL 的 algorithm 和 protocol 是两个正交维度： algorithm: rank 之间怎样组织数据流 Ring / Tree / CollNet / NVLS ... protocol: 一条逻辑边上怎样编码 payload、发布 ready 状态并推进 FIFO step LL / LL128 / Sim...

本章问题 “Tree 小消息快，Ring 大消息快”只能当作待验证的经验，不能当作 NCCL 原理。它省略了至少五个变量： Tree 是平衡二叉树，还是节点内链 + 节点间树？ Tree 和 Ring 是否使用同一个 protocol？ 消息被多少个 channel 切分？ world size 和 node 数分别是多少？ 当前拓扑搜索究竟生成了什么 parent/children？ ...

Ring AllReduce 不只是“绕一圈” 最常见的描述： Ring AllReduce = ReduceScatter + AllGather 这句话是正确的，但不足以读源码或排障。 还需要回答： chunk 编号按 user rank，还是 ring position？ 每个 rank 第一次发送哪个 chunk？ recvReduceSend 连续执行多少...

面对 20% 回归，先别调环境变量 一条 AllReduce 变慢，可能来自： 应用层: collective sequence / bucket / overlap / straggler enqueue/tuning: algorithm / protocol / channel / chunk device: kernel occupancy / contentio...

一条直线为什么解释不了 1 B 到 1 GiB 通信性能最常见的模型： [T(n) = alpha + n / beta] 其中： alpha：与 payload 无关的固定时间； beta：稳态有效 payload bandwidth； n：消息字节数； T：完成时间。 本机对 1 B 到 1 GiB 的 31 个 AllReduce median 做一条全局...

“跑十次取平均”为什么不够 同一台 4×V100 机器、同一个 1 MiB AllReduce，本章实际测到： -I 0: 50.385 us -I 1: 54.125 us -z 0: 60.730 us -z 1: 77.315 us -z 2: 77.905 us -z 3: 88.210 us -n 1:...

122 GB/s 到底是什么 在四张 V100 上运行 all_reduce_perf，经常会看到类似结果： size time algbw busbw 256 MiB 3290 us 81.59 GB/s 122.38 GB/s 但第 2 章实测单个 GPU pair 的单向 P2P 吞吐约为 48.3 GB/s。 如果把 122...

“AllReduce = ReduceScatter + AllGather”少了什么 这句等式在分布式训练中经常出现： [\text{AllReduce} \text{ReduceScatter} + \text{AllGather}] 它至少有四种不同含义： 输出语义等价：每个 rank 最终拿到相同完整 tensor。 通信步数等价：Ring 模型中两段的 step ...

本章要区分的四件事 下面四句话经常被混在一起： 所有 rank 必须按相同顺序调用 collective。 ncclGroupStart/End 可以批量提交多次 NCCL 调用。 ncclGroupEnd 返回时，操作已经进入 CUDA stream。 CUDA stream 上的通信已经完成。 只有前两句都正确，第三句也不推出第四句。 本章回答： “相...

本章不是 API 名词表 知道 AllReduce 是“归约后广播”还不够。线上问题通常出现在更具体的地方： count 对 AllGather 表示单个 rank 的输入元素数，还是总接收元素数？ ReduceScatter 的输入为什么必须有 nranks * recvcount 个元素？ AllGather 和 ReduceScatter 的 in-place 指针为...

本章要解决什么问题 下面三种代码看起来分别是同步、异步后等待和显式设备同步： # A dist.all_reduce(x, async_op=False) # B work = dist.all_reduce(x, async_op=True) work.wait() # C work = dist.all_reduce(x, async_op=True) torch.cuda.sy...

本章要解决什么问题 排查 NCCL 问题时，经常会看到下面几条信息： nvidia-smi: CUDA Version 13.0 nvcc: release 12.6 torch.version.cuda: 12.6 torch.cuda.nccl.version(): 2.22.3 dpkg: libnccl2 2.22.3-1+cuda12.6 NCCL INFO: NCCL ver...

本章要解决什么问题 一条 NCCL 初始化日志可能是： ncclCommInitRank comm 0x... rank 2 nranks 4 cudaDev 2 nvmlDev 2 busId 1d000 ... Init COMPLETE 这里至少有五种身份容易混在一起： global rank process-group rank NCCL communicator rank...

本章要解决什么问题 nvidia-smi topo -m 在本机输出： GPU0 GPU1 GPU2 GPU3 CPU Affinity NUMA Affinity GPU0 X NV2 NV2 NV2 0,2,4,... 0 GPU1 NV2 X NV2 NV2 0,2,4,... 0 GPU...

这个专题是一套 36 章 NCCL 专家课程。每章都按“提出可证伪假设、直接阅读固定版本源码、运行控制变量实验、统计结果、用源码解释结果、明确结论边界”的顺序完成，不再把环境检查或文件路径列表当成教程。 实验环境的核心信息： GPU: 4 x Tesla V100-SXM2-32GB GPU 拓扑: GPU 两两 NV2 CUDA Toolkit: 12.6 PyTorch: 2.5.0...