2、NCCL核心概念:通信域与集合通信操作
好,咱们今天聊聊NCCL里最核心的几个概念。说实话,我刚接触NCCL的时候,也被这些术语搞得有点晕。但后来我发现,只要搞懂了通信域和那几种集合通信操作,基本上就能上手干活了。
2.1 通信域(Comm)—— 你的“通信圈子”
通信域,英文叫 Communicator,简称 Comm。你可以把它理解成一个“通信圈子”。
举个例子。你有一群 GPU,它们要一起干活。但你不能让所有 GPU 都乱哄哄地互相喊话吧?那不乱套了?所以 NCCL 搞了个“通信域”的概念。你把需要互相通信的 GPU 拉到一个组里,这个组就是一个 Comm。
我个人习惯把 Comm 想象成一个群。群里的人(GPU)可以发消息(数据),群主(NCCL)负责管理。每个 GPU 在群里都有一个唯一的编号,叫 rank。比如 rank 0、rank 1、rank 2……
关键点来了:
- 每个进程只能属于一个 Comm 吗? 不是。一个进程可以加入多个 Comm,就像你可以同时加好几个群。
- Comm 怎么创建? 通常用
ncclCommInitRank或者ncclCommInitAll。我建议你尽量用ncclCommInitRank,因为它更灵活,能让你控制每个 rank 的初始化顺序。
避坑指南: 我曾经在项目里犯过一个低级错误——多个进程用了同一个 Comm 的 rank 0。结果通信直接乱掉,数据对不上。后来排查了半天才发现,原来是我初始化时传错了参数。记住:每个进程的 rank 必须唯一,且连续。
2.2 集合通信操作 —— 核心的四种“招式”
搞定了通信域,接下来就是怎么在域里通信了。NCCL 提供了几种集合通信操作,说白了就是“大家一起干一件事”。我挑最常用的四种来讲。
2.2.1 AllReduce —— 最常用的“求和”
AllReduce 是分布式训练里最常用的操作。它的作用:每个 GPU 上有一份数据,大家把自己的数据加起来,然后每个 GPU 都拿到最终的和。
你想想看,在数据并行训练中,每个 GPU 算完梯度后,是不是要把梯度加起来?对,就是用 AllReduce。
代码示例:
// 假设每个 GPU 上有 1024 个 float 数据
float* sendbuf; // 本地数据
float* recvbuf; // 接收结果
ncclAllReduce(sendbuf, recvbuf, 1024, ncclFloat, ncclSum, comm, stream);
这里 ncclSum 是操作类型,你也可以用 ncclProd(乘积)或 ncclMax(最大值)。但我个人经验,99% 的场景下你只会用到 ncclSum。
2.2.2 AllGather —— “收集”所有人的数据
AllGather 的意思是:每个 GPU 把自己的一块数据拿出来,然后大家互相交换,最后每个 GPU 都拥有所有人的完整数据。
举个例子。你有 4 个 GPU,每个 GPU 上有 256 个数据。AllGather 之后,每个 GPU 上都会有 1024 个数据(4 * 256)。
这个操作在模型并行中很常见。比如你要做张量并行,需要把切分后的参数收集回来。
小技巧: 如果你只需要把数据收集到一个 GPU 上(比如 rank 0),那就用 ncclGather,而不是 AllGather。AllGather 是“所有人都有”,Gather 是“一个人有”。别搞混了。
2.2.3 ReduceScatter —— “先求和,再分发”
ReduceScatter 是 AllReduce 的“变种”。它先把数据做规约(比如求和),然后把结果分散到各个 GPU 上。
我刚开始学的时候,总觉得 ReduceScatter 和 AllReduce 差不多。后来才明白:AllReduce 是“每个人都拿到完整结果”,ReduceScatter 是“每个人只拿到一部分结果”。
为什么要有这个操作?因为有些场景下,你不需要每个人都拿到完整数据。比如在做梯度累积时,你可以用 ReduceScatter 把梯度分散到不同 GPU 上,减少通信量。
代码示例:
// 每个 GPU 有 1024 个数据,4 个 GPU
// 结果:每个 GPU 得到 256 个数据(1024 / 4)
ncclReduceScatter(sendbuf, recvbuf, 256, ncclFloat, ncclSum, comm, stream);
2.2.4 Broadcast —— “一个人说话,所有人听”
Broadcast 最简单:一个 GPU(通常是 rank 0)把自己的数据发给所有其他 GPU。
这个操作常用于初始化参数。比如你加载了一个预训练模型,只需要 rank 0 读一次,然后 Broadcast 给所有其他 GPU。
注意: Broadcast 是“一对多”,不是“多对多”。如果你想让所有 GPU 都互相发数据,那得用 AllGather 或者 AllReduce。我曾经见过有人用 Broadcast 循环发数据,结果性能惨不忍睹。
2.3 同步与异步 —— 别让 GPU 等你
NCCL 的操作可以是同步的,也可以是异步的。默认情况下,NCCL 操作是异步的——你调用 ncclAllReduce 后,函数会立即返回,不会等操作完成。
为什么要异步?因为 GPU 计算和通信可以重叠。你可以在 GPU 做计算的同时,让 NCCL 在后台传输数据。这样能最大化利用硬件资源。
但异步也有坑。如果你在通信还没完成时就读取结果,那数据肯定是错的。
怎么保证同步? 两种方式:
- 用 CUDA Stream 同步:
cudaStreamSynchronize(stream)。这会阻塞 CPU,直到 stream 上的所有操作完成。 - 用 NCCL 的
ncclGroupStart/End: 把多个 NCCL 操作放到一个组里,然后一起同步。
我个人习惯用第一种。简单直接,不容易出错。
避坑指南: 我曾经在调试一个分布式训练任务时,发现 loss 死活不收敛。查了两天,最后发现是 NCCL 操作没同步——我在通信还没完成时就开始用数据了。加上 cudaStreamSynchronize 后,问题立刻解决。嗯,这里要注意:异步虽好,但别贪杯。
2.4 知识体系图
下面这张图总结了本章的核心内容。你可以把它当作一个“速查表”。
2.5 总结
好了,这一章的内容就这些。总结一下:
- 通信域是 GPU 的“群”,rank 是群里的唯一 ID。
- 四种集合通信操作:AllReduce(求和)、AllGather(收集)、ReduceScatter(求和+分发)、Broadcast(广播)。
- 异步是默认行为,记得用
cudaStreamSynchronize同步。
这些概念是 NCCL 的基石。搞懂了它们,后面讲性能优化、故障排查时,你才能跟得上。嗯,下一章咱们聊聊 NCCL 的初始化流程和常见配置,到时候见。
个人建议: 如果你刚开始学 NCCL,别急着看高级特性。先把这四种操作练熟。我当年就是太心急,结果基础没打好,后面调试起来特别痛苦。