📖 阅读时间:约 8 分钟
前两篇,我们先把集群通信的两层基本逻辑铺开了。
连载 01 讲 All-Reduce:它不是把所有梯度先收拢到一个中心节点,而是通过 Reduce-Scatter 和 All-Gather,让多张 GPU 一边归约、一边交换,避免参数服务器成为单点瓶颈。
连载 02 讲 Tree:通信算法不只有 Ring。Tree 用对数级的层次传播减少通信轮次,尤其适合启动延迟占主导的小消息;Ring 则更擅长让大消息沿链路持续流动。
但知道“数据该怎么绕一圈”之后,还有一个更贴近性能现场的问题:算法路径已经选对,为什么 GPU 间链路的实际使用率还是上不去?
一条 1 GB 的梯度要穿过 GPU 间链路,理论带宽可能很高,监控里却经常出现一段忙、一段空。原因不一定是链路变窄了,而可能是整块数据等待就绪、每一步通信串行启动、归约和转发互相等候,导致宝贵的传输窗口被空档切碎。
如果等 1 GB 全部准备好再一口气搬,链路会在“等待—突发传输—再次等待”之间反复切换。真正需要优化的,不只是总共搬了多少字节,而是能不能让不同数据块连续进入不同处理阶段,让链路尽量保持有货可运。
本文真正要回答的是:物理链路没变,NCCL 如何用 Channel、Chunk 和 Slice 把低利用率的间歇搬运,改造成持续推进的通信流水线?
NCCL 的做法可以先压缩成一句话:Channel 提供逻辑并行车道,Chunk 决定每辆车装多少货,Slice 让车内的小批货也能细粒度交接。

📋 本文目录
一、What:Channel、Chunk、Slice 到底是什么?
1. Channel:一次 collective 内的逻辑并行通道
2. Chunk:沿算法步次流动的数据块
3. Slice:协议缓冲区内更细的交接单元
二、Why:为什么不整块传,而要切成流水线?
三、How:这条流水线是怎样跑起来的?
1. CTA 为活跃 Channel 提供执行班组
2. Protocol 规定“怎么装货、怎么确认到站”
3. Topology 决定逻辑流水线铺在哪里
四、效果小结:路没变宽,运输方式变了
一、What:Channel、Chunk、Slice 到底是什么?
先把三个常被混用的概念分层。
1. Channel:一次 collective 内的逻辑并行通道
Channel 不是一根新增的 NVLink,也不是一条独占网线。它是 NCCL 在一次 collective 中划分出来的逻辑工作分区:每个 Channel 持有自己的通信拓扑,例如一条 Ring 或一棵 Tree,并处理输入数据的一部分。
因此,把 Channel 从 1 增加到 8,不等于物理带宽变成 8 倍。多个 Channel 可能仍共享同一条物理链路;增加的是可独立推进的工作和在途数据,而不是凭空复制硬件。
2. Chunk:沿算法步次流动的数据块
Chunk 是 collective 算法真正向前推进的数据单元。
以 Ring AllReduce 为例,张量先分配到不同 Channel;每个 Channel 再按 rank 和轮次取出 Chunk。一个 Chunk 在 Reduce-Scatter 阶段逐跳传输并归约,得到最终结果后,再在 All-Gather 阶段逐跳分发。NCCL 官方文档对 AllReduce 的语义定义是“先归约,再把结果提供给所有 rank”,而 Chunk 正是这套算法执行时可持续流动的数据块(参考 NVIDIA NCCL Collectives)。

3. Slice:协议缓冲区内更细的交接单元
Slice 比 Chunk 更细。一个 Chunk 不必等到整体完成后才交给下一跳,它可以继续切成 Slice,在协议缓冲区中分批 load、reduce、send、recv 和确认。
可以把三者记成一条层级:
Channel 决定“有几条逻辑车道”,Chunk 决定“每辆车装多少”,Slice 决定“车内货物以多细的批次交接”。
二、Why:为什么不整块传,而要切成流水线?
先看最朴素的串行方案:一整块数据完成传输后,下一环节才开始归约;归约完成后,下一跳才继续转发。
这种方案的问题不是链路绝对慢,而是阶段之间存在大量等待。传输工作时,部分计算资源可能空闲;归约工作时,链路又可能没有新数据可发。一次 collective 经过多跳、多轮后,这些空隙会不断累积。
把大块拆成多个 Chunk 后,时间线会改变:
Chunk A 已到下一站,开始归约; Chunk B 正在链路上传输; Chunk C 还在本站装载; 更早完成的 Chunk 可能已经进入下一轮转发。
于是传输、归约和转发不再严格串行,而是在不同 Chunk 上重叠。收益来自等待被折叠、路径利用率提高,不来自物理带宽被放大。

但 Chunk 也不是越小越好。块太大,流水线启动慢;块太小,循环、同步、flag 和协议控制开销占比会上升。真正目标是找到足够细、又不会被控制开销吞掉的推进粒度。
三、How:这条流水线是怎样跑起来的?
1. CTA 为活跃 Channel 提供执行班组
在常见 NCCL collective kernel 中,可以把 CTA 理解为一个 CUDA block。当前主干实现会通过 channelMask 把 blockIdx.x 映射到启用的 Channel:一个 block 领取一个 Channel 的工作,block 内线程协作完成 load、reduce、copy、send 和 recv(见 NCCL 源码 common.h)。
这是一种理解模型,不是跨版本不变的硬规则。稳定成立的是:Channel 是调度与拓扑上的工作分区,CTA 和线程则是把这些工作真正执行出来的 GPU 资源。
2. Protocol 规定“怎么装货、怎么确认到站”
NCCL 的 Protocol 决定数据如何放进通信缓冲区、发送端与接收端如何同步,以及每次循环能搬多少有效数据。
- Simple
更偏向大消息带宽; - LL
用更细的控制粒度换低延迟; - LL128
在低延迟与带宽之间折中。
协议会影响 Slice 粒度、缓冲区占用和同步方式,因此“Channel 数”“Chunk 大小”和“Protocol”不能分开看。NCCL 提供 NCCL_PROTO、NCCL_BUFFSIZE、NCCL_MIN_CTAS、NCCL_MAX_CTAS 等环境变量用于实验和调优,但官方也提示,调试类变量不宜在生产环境长期固定(参考 NCCL Environment Variables)。

3. Topology 决定逻辑流水线铺在哪里
NCCL 会根据 GPU、NVLink、PCIe、CPU、NIC 和交换结构选择 Ring、Tree 等算法及通信路径。Channel 在这些路径上建立逻辑并行,但多个 Channel 仍可能争用同一物理链路。
这解释了为什么 Channel 不是越多越好:
太少,可能没有足够工作把链路喂满; 适中,可以让更多 Chunk 保持在途并覆盖阶段等待; 太多,会增加 kernel 启动、同步、buffer、寄存器、共享内存和 SM 压力,还可能放大共享链路争用。

四、效果小结:路没变宽,运输方式变了
NCCL 的 Channel/Chunk 流水线可以压缩成四步:
- 按拓扑选路线
:决定 Ring、Tree 和实际通信路径; - 用 Channel 分区
:让一次 collective 拥有多条可独立推进的逻辑工作线; - 用 Chunk 持续推进
:让不同算法步骤在不同数据块上重叠; - 用 Slice 与 Protocol 交接
:让缓冲区内部也能细粒度流动。
最终收益是更高的路径利用率和更少的阶段空等。边界也同样清楚:Channel 太少可能喂不满链路,太多会增加资源和同步压力;Chunk 太大流水线启动慢,太小则控制开销占比上升。
所以,真正值得记住的不是某个“神奇 Channel 数”,而是这条调优原则:
先让 NCCL 根据拓扑自动选择,再用可复现实测观察链路利用率、延迟与 SM 占用;只有默认选择确实失配时,才收紧 CTA、协议或缓冲区边界。
「老许漫谈AIInfra」 · 持续关注 AI 基础设施工程实践
夜雨聆风