乐于分享
好东西不私藏

【AI集群通信连载 03】NCCL Channel/Chunk 如何把链路变成流水线?

【AI集群通信连载 03】NCCL Channel/Chunk 如何把链路变成流水线?

📖 阅读时间:约 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)。

Ring AllReduce 中 Chunk 的两段旅程:先归约,再分发

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 流水的差异:物理链路未变,阶段空等显著减少

但 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_PROTONCCL_BUFFSIZENCCL_MIN_CTASNCCL_MAX_CTAS 等环境变量用于实验和调优,但官方也提示,调试类变量不宜在生产环境长期固定(参考 NCCL Environment Variables)。

NCCL 三层推进粒度:Channel、Chunk、Slice,以及 CTA 与 Protocol 的位置

3. Topology 决定逻辑流水线铺在哪里

NCCL 会根据 GPU、NVLink、PCIe、CPU、NIC 和交换结构选择 Ring、Tree 等算法及通信路径。Channel 在这些路径上建立逻辑并行,但多个 Channel 仍可能争用同一物理链路。

这解释了为什么 Channel 不是越多越好:

  • 太少,可能没有足够工作把链路喂满;
  • 适中,可以让更多 Chunk 保持在途并覆盖阶段等待;
  • 太多,会增加 kernel 启动、同步、buffer、寄存器、共享内存和 SM 压力,还可能放大共享链路争用。
Channel 调优边界:过少喂不满,过多引入资源与同步争用

四、效果小结:路没变宽,运输方式变了

NCCL 的 Channel/Chunk 流水线可以压缩成四步:

  1. 按拓扑选路线
    :决定 Ring、Tree 和实际通信路径;
  2. 用 Channel 分区
    :让一次 collective 拥有多条可独立推进的逻辑工作线;
  3. 用 Chunk 持续推进
    :让不同算法步骤在不同数据块上重叠;
  4. 用 Slice 与 Protocol 交接
    :让缓冲区内部也能细粒度流动。

最终收益是更高的路径利用率和更少的阶段空等。边界也同样清楚:Channel 太少可能喂不满链路,太多会增加资源和同步压力;Chunk 太大流水线启动慢,太小则控制开销占比上升。

所以,真正值得记住的不是某个“神奇 Channel 数”,而是这条调优原则:

先让 NCCL 根据拓扑自动选择,再用可复现实测观察链路利用率、延迟与 SM 占用;只有默认选择确实失配时,才收紧 CTA、协议或缓冲区边界。

「老许漫谈AIInfra」 · 持续关注 AI 基础设施工程实践