夜雨聆风Ring AllReduce 物理数据路径
聚焦一个问题：一个数据块在 Ring AllReduce 中，机内怎么传、跨机怎么传，GPU 和 IB 卡各扮演什么角色。 全部结论附 NCCL 源码出处，可自行核对。
1. Ring 的节点是 GPU，不是网卡
结论
Ring = GPU（rank）首尾相连组成的逻辑环。IB 网卡不是环上的节点，只是环上"跨节点那一段"的物理管道。
源码证据
struct ncclRing 存的是 rank 号，不是网卡号：
nccl/src/include/device.h:175
struct ncclRing {
int prev; // 左邻居的 rank 号
int next; // 右邻居的 rank 号
...
};
ring 数组遍历时走的是 rank：
nccl/src/graph/rings.cc:42-46
int current = rank;
for (int i = 0; i < nranks; i++) {
rings[r * nranks + i] = current; // 存的全是 rank
current = next[r * nranks + current]; // 沿 next 跳到下一个 rank
}
每个 channel 内嵌 1 个 ring，ring 编号 = channel 编号：
nccl/src/include/comm.h:155
struct ncclChannel {
struct ncclRing ring; // 每 channel 1 个 ring
...
};
nccl/src/graph/rings.cc:48
snprintf(prefix, "Channel %02d/%02d :", r, nrings); // ring 编号直接打印成 channel 编号
2. 环数 = channel 数，1:1
结论
源码里 nrings 直接传的就是 nChannels。你设 64 channel 就是 64 个环。MAXCHANNELS 硬上限 64。
源码证据
nccl/src/graph/connect.cc:522
ncclBuildRings(nChannels, rings, comm->rank, comm->nRanks, ringPrev, ringNext);
// ^^^^^^^^^^ nrings 直接传 nChannels
nccl/src/include/device.h:91
#define MAXCHANNELS 64 // 硬上限
nccl/src/graph/connect.cc:504-505
nChannels = copyChannels(comm, nChannels, std::max(ncclMinNchannels(), comm->config.minCTAs), ...);
// 用 copyChannels 把基础拓扑搜出来的 channel 复制到你设的数量
重要细节：64 个环里很多是复制品
16 卡拓扑搜索可能只产出少数几种不同的 ring 排列，剩下的靠 copyChannels 复制。所以 64 个环里大量"逻辑拓扑完全相同"的环——同一个 rank 在这些环里 prev/next 一模一样。区别只在物理层：64 个 channel 分摊到 NIC 上（比如 8 张 NIC，每张扛 8 个 channel 的 QP）。
所以之前文档里"1个环×8个Channel"的措辞是错的，准确说法：
8个 channel 各持 1 个 ring，这 8 个 ring 的 GPU 顺序相同，区别只在跨节点段走哪张 IB 卡。
3. QP 是网卡的，不是 GPU 的
结论
QP 建在 IB 网卡上。GPU 只是数据的源头和目的地，不持有 QP。
跨节点段的完整链路：
源GPU显存 → PCIe → 本机IB网卡(这里才有QP) → 线缆 → 对端IB网卡(QP) → PCIe → 对端GPU显存
GPU 之间不"通过 QP 交互"，是两块网卡的 QP 之间交互，GPU 只负责出/收数据。
补充：机内段根本没有 QP
机内相邻 GPU 之间走 NVLink（经 NVSwitch），不经过 IB 网卡，没有 QP，没有封包。NCCL 用 P2P transport（不是 net transport）处理机内段。
4. 一个 chunk 绕环一圈的物理路径
场景
3 节点 × 2 GPU，环序 0→1→2→3→4→5→0
节点A: GPU0, GPU1
节点B: GPU2, GPU3
节点C: GPU4, GPU5
环序: GPU0 → GPU1 → GPU2 → GPU3 → GPU4 → GPU5 → GPU0
|----机内----|----跨机----|----机内----|----跨机----|----机内----|----跨机----|
6 跳逐跳拆解
跳1 GPU0→GPU1 机内（NVLink） GPU0显存 → NVSwitch → GPU1显存 不经过 PCIe到CPU，不经过 IB 卡，没有 QP，没有封包。 带宽 300GB/s 级别。
跳2 GPU1→GPU2 跨机（IB） GPU1显存 → PCIe → 节点A的IB网卡 → 线缆 → 节点B的IB网卡 → PCIe → GPU2显存 这段才有 QP，才有 RDMA_WRITE，才有封包/解包。 带宽受限于 IB 卡 400G = 50GB/s。
跳3 GPU2→GPU3 机内（NVLink） GPU2显存 → NVSwitch → GPU3显存
跳4 GPU3→GPU4 跨机（IB） GPU3显存 → PCIe → 节点B的IB网卡 → 线缆 → 节点C的IB网卡 → PCIe → GPU4显存
跳5 GPU4→GPU5 机内（NVLink） GPU4显存 → NVSwitch → GPU5显存
跳6 GPU5→GPU0 跨机（IB） GPU5显存 → PCIe → 节点C的IB网卡 → 线缆 → 节点A的IB网卡 → PCIe → GPU0显存
统计
6 跳 = 3 跳机内 + 3 跳跨机。机内那 3 跳基本不花时间（NVLink 又宽又快），真正慢的是跨机那 3 跳——每跳都要走 PCIe 上 IB 卡、封包、过线缆、解包、再 PCIe 下到显存。
5. 为什么跨机段是瓶颈
带宽对比
跨机段的 IB 卡 50GB/s 是整条链路最窄的地方，比 NVLink 慢 6 倍，比显存带宽慢近 100 倍。
Ring AllReduce 的优化全盯着跨机段
机内段白送，跨机段是瓶颈。8 卡 8 NIC 的设计就是让跨机那 3 跳分散到不同 IB 卡上并行走，而不是挤一张卡。
8 个 channel（8 个 ring，拓扑相同）跨节点段分别经 8 张 IB 卡中转，8×400G = 3.2Tbps 打满。这就是你 ring-allreduce-logic-vs-physics.html 下图画的：同一个环逻辑序，跨节点时数据经 GPU0→IB0、GPU1→IB1...GPU7→IB7 各走各的卡。
6. 跨机段 RDMA 传一个 chunk 的完整动作
以 GPU1 给 GPU2 发一个 chunk 为例（Reduce-Scatter 的一步）：
发送端（节点A 的 GPU1）
- NCCL 填 WQE（任务单）：
操作码 = RDMA_WRITE - 源地址 = GPU1 显存里这个 chunk 的地址（IOVA）
- 源 lkey = 节点A MR 的本地 key
- 目标远程地址 = GPU2 显存里对应位置（节点B 的 IOVA）
- 目标 rkey = 节点B MR 产生的远程 key
- 长度 = chunk 大小
NCCL 调 ibv_post_send() → 写 Doorbell（mmio 直写网卡寄存器，0 次系统调用）
- 节点A 的 IB 网卡硬件执行（CPU 不参与）：
DMA 读 GPU1 显存（经 PCIe，用 lkey 验证权限，IOMMU 翻译地址） - 封包：加 IB 头（目标 LID/GID + QPN + PSN），RoCEv2 模式再加 UDP/IP/以太网头
- 从 400G 物理口发出
网络传输
IB 模式：IB 交换机查 LID 路由表转发 RoCEv2 模式：以太网交换机查 MAC/IP 转发，PFC/ECN 保证不丢包
接收端（节点B 的 GPU2）
- 节点B 的 IB 网卡硬件执行（CPU 不参与）：
收包，解析包头 - 提取 QPN → 查 QP 表 → 找到对应 QP
- 提取 rkey → 查 MR 表 → 匹配 MR
- IOMMU 验证 rkey → 允许
- DMA 写数据到 GPU2 显存（经 PCIe，IOVA→IOMMU→物理页帧）
网卡 DMA 写 CQE 到 CQ → NCCL 轮询 ibv_poll_cq() 得到"收完了"
CPU 参与度
全程 CPU 只做两件事：
- 初始化时（建 QP、注册 MR、交换连接信息）
- 运行时轮询 CQ（知道发完了/收完了）
数据搬运（显存→PCIe→网卡→线缆→网卡→PCIe→显存）全是硬件完成的。
7. 8 卡 8 NIC 场景的 channel→NIC 映射
crossNic=0（默认，你的场景）
每个 channel 的跨机段固定走 1 张 IB 卡，不同 channel 走不同 IB 卡：
channel-0 (ring-0) 跨节点段 → GPU0 → IB0 → 对端IB0 → 对端GPU0
channel-1 (ring-1) 跨节点段 → GPU1 → IB1 → 对端IB1 → 对端GPU1
...
channel-7 (ring-7) 跨节点段 → GPU7 → IB7 → 对端IB7 → 对端GPU7
如果 channel 数 > NIC 数（比如 64 channel、8 NIC）： channel-07 → IB07（第一轮） channel-815 → IB07（第二轮） ... 多个 channel 共享同一张 IB 卡，QP 也在这张卡上排队。
源码出处
nccl/src/graph/search.cc:827 (注释)
Ring: NET n -> GPU a -> GPU b -> .. -> GPU x -> NET n (or m if crossNic)
nccl/src/graph/search.cc:16
NCCL_PARAM(CrossNic, "CROSS_NIC", 2); // 默认值 2（自动）
crossNic=0：每个 channel 固定走 1 张 NIC crossNic=1：允许同一 channel 跨不同 NIC
8. 常见误解纠偏表
9. Ring 的意义不是为了"拿到数据"，是用最少连接把带宽吃满
"拿到所有数据"不是 ring 被选中的原因
全互联(mesh) 也能拿到所有数据，broadcast 也能拿到。ring 不是唯一能做到的方式。
ring 解决的是两个约束同时满足：
约束1：连接数要少 1024 卡全互联，每台机器连 1023 个邻居，QP 爆炸。 ring 每个 GPU 只连 prev/next 两个邻居，QP 数量可控。
约束2：带宽要打满 如果只有 1 个 channel（1 个 ring），跨机段 1 张 IB 卡 400G，只吃 1/8 带宽。 8 个 channel（8 个 ring 拓扑相同）跨节点段分别走 8 张 IB 卡，8×400G=3.2Tbps 全吃满。
ring 的设计让每跳只传 1/N 的数据量（N 是 GPU 数），所有 GPU 同时传，流水线并行。这样跨机段那个 50GB/s 的瓶颈被多个 channel 叠加 + 流水线填补，尽量逼近 NVLink 的速度。
一句话：ring 不是为了"拿到数据"，是为了"用最少的连接数（每GPU只连2邻居）把跨机带宽吃满"。拿到数据是结果，不是 ring 被选中的原因——mesh 也能拿到数据，但 QP 数量会炸。
10. 单个 chunk 一次传输的完整例子
场景设定
4 GPU，2 节点，每节点 2 GPU：
节点A：GPU0 (rank0)、GPU1 (rank1)
节点B：GPU2 (rank2)、GPU3 (rank3)
环序：GPU0 → GPU1 → GPU2 → GPU3 → GPU0
机内 跨机 机内 跨机
数据：每个 GPU 有一组梯度，切成 4 块（chunk0~chunk3）。 追踪：GPU0 的 chunk0，256MB，要沿环传到 GPU1。
chunk0 在 GPU0 显存的地址：0x70000000
chunk0 大小：256MB = 268435456 字节
这是 Reduce-Scatter 阶段的第 1 步：GPU0 把 chunk0 发给 next（GPU1）。
跳1：GPU0 → GPU1（机内，NVLink）
GPU0 和 GPU1 同在节点A，通过 NVSwitch 互联。
谁在干活
NCCL 通道(channel-0) 的 P2P transport（不是 net transport）
没有 QP
没有 IB 网卡
没有 WQE
没有封包
没有 CPU 参与
物理路径
GPU0显存(0x70000000)
↓
NVSwitch（节点A内部）
↓
GPU1显存(0x71000000)
就这 3 步。GPU0 的 DMA 引擎直接通过 NVLink 把数据写到 GPU1 显存，GPU1 的 CPU 完全不知道。
具体动作
- NCCL 在 GPU0 上 launch 一个 CUDA kernel
- kernel 执行：从地址 0x70000000 读 256MB，通过 NVLink 写到 GPU1 的 0x71000000
- NVSwitch 负责转发（节点A内 8 卡全互联，NVSwitch 知道 GPU1 在哪）
- GPU1 显存 0x71000000 处出现 256MB 数据
耗时估算
256MB ÷ 300GB/s ≈ 0.85μs
GPU1 收到后做什么
GPU1 把收到的 chunk0（GPU0的数据）加到自己对应的 chunk0 上：
GPU1.chunk0 = GPU1.chunk0 + 收到的chunk0
= [2,2,2,2] + [1,1,1,1]
= [3,3,3,3]
加完之后，这个 [3,3,3,3] 就是下一步要发给 GPU2 的数据。
跳2：GPU1 → GPU2（跨机，IB 网卡 + RDMA）
GPU1 在节点A，GPU2 在节点B，必须走 IB 网卡。
谁在干活
NCCL 通道(channel-0) 的 net transport
节点A 的 IB 卡（CX7-1，挂在 GPU1 旁边）
QP：QP-100（发）在节点A IB卡上，QP-200（收）在节点B IB卡上
MR：节点A 注册了 GPU1 显存，节点B 注册了 GPU2 显存
CPU 参与：只填 WQE + 敲 Doorbell + 轮询 CQ
初始化时已经建好的东西（只做一次）
节点A：
MR-A 注册了 GPU1 显存 0x71000000，长度 1GB
lkey = 0x1A2B（节点A自己用）
rkey = 0x3C4D（给对端用，但本场景是RDMA_WRITE所以对端不需要A的rkey）
IOVA = 0x10000000（网卡视角的地址，和CPU虚拟地址0x71000000不同）
QP-100（发QP）建在节点A的CX7-1上
寄存器里写死了对端信息：
远程LID = 12（节点B CX7-1的LID）
远程QPN = 200（节点B的收QP号）
远程PSN = 0（起始序列号）
节点B：
MR-B 注册了 GPU2 显存 0x72000000，长度 1GB
lkey = 0x5E6F（节点B自己用）
rkey = 0xABCD（给节点A用，A写B时填这个）
IOVA = 0x20000000
QP-200（收QP）建在节点B的CX7-2上
寄存器里写死了对端信息：
远程LID = 5（节点A CX7-1的LID）
远程QPN = 100
远程PSN = 0
NCCL 填写 WQE（任务单）
NCCL 在节点A 填一张 WQE，塞进 QP-100 的 SQ（发送队列）：
WQE 内容：
操作码： RDMA_WRITE
源地址(IOVA)： 0x10000000 ← GPU1显存的IOVA（不是CPU虚拟地址0x71000000）
源lkey： 0x1A2B ← 节点A MR-A的本地key
长度： 268435456 ← 256MB
目标远程地址： 0x20000000 ← GPU2显存的IOVA（不是CPU虚拟地址0x72000000）
目标rkey： 0xABCD ← 节点B MR-B产生的远程key（不是A的）
关键：目标 rkey 是 B 的 MR 产生的。A 要写 B 的显存，得拿 B 给的钥匙。
NCCL 敲 Doorbell
NCCL 往 Doorbell 寄存器写 1 个值（mmio 直写网卡硬件，0 次系统调用，纳秒级）
→ 节点A 的 CX7-1 网卡被通知："QP-100 有新任务"
节点A 网卡硬件执行（CPU 不参与）
第1步：网卡 DMA 读 WQE
CX7-1 通过 PCIe 读 QP-100 的 SQ，拿到刚才填的 WQE
第2步：网卡 DMA 读 GPU1 显存
WQE 说源地址=0x10000000，长度=256MB
网卡通过 PCIe → IOMMU 翻译 IOVA 0x10000000 → 物理页帧 → 读 GPU1 显存
用 lkey=0x1A2B 验证权限：IOMMU 查 MR 表，0x1A2B 对应 MR-A，允许读
256MB 数据被 DMA 读进网卡内部的发送缓冲区
第3步：网卡封包
把 256MB 切成多个网络包（每个包 payload 约 4KB，约 65536 个包）
每个包头加：
目标LID = 12 ← IB交换机靠这个转发
目标QPN = 200 ← 到了节点B后网卡靠这个找到QP-200
PSN = 0,1,2,... ← 每个包+1，用于检错和硬件重传
源LID = 5
源QPN = 100
如果是 RoCEv2 模式，还要加 UDP头(端口4791) + IP头 + 以太网头
第4步：网卡从 400G 物理口发出
65536 个包依次从 CX7-1 的 400G 端口发出去
网络传输
节点A CX7-1 → IB线缆 → IB交换机 → IB线缆 → 节点B CX7-2
IB交换机查路由表：目标LID=12 → 从端口7转发
（交换机只看LID，不看QPN/PSN，QPN/PSN是端点级别，交换机不管）
节点B 网卡硬件执行（CPU 不参与）
第5步：网卡收包，解析包头
CX7-2 收到包，解析出：
目标QPN = 200 → 查 QP 表 → 找到 QP-200
PSN = 0 → QP-200 期望的起始PSN=0 ✓ → 接收
PSN = 1 → 期望1 ✓ → 接收
...（如果某个PSN不连续 → 发NACK → 节点A网卡硬件重传，微秒级）
第6步：网卡查 MR 表
包头里有 rkey = 0xABCD
CX7-2 查 MR 表：0xABCD → MR-B → 目标地址 0x20000000
IOMMU 验证 rkey → 允许写
第7步：网卡 DMA 写 GPU2 显存
WQE 说目标地址=0x20000000（IOVA），长度=256MB
CX7-2 通过 PCIe → IOMMU 翻译 IOVA 0x20000000 → 物理页帧 → 写 GPU2 显存
256MB 数据被 DMA 写进 GPU2 显存 0x72000000（CPU虚拟地址）
第8步：网卡写 CQE
CX7-2 DMA 写一个 CQE 到节点B 的 CQ
→ 节点B 的 NCCL 轮询 ibv_poll_cq() → 知道"收完了"
同时，节点A 的完成通知
CX7-1 发完所有包后，DMA 写一个 CQE 到节点A 的 CQ
→ 节点A 的 NCCL 轮询 ibv_poll_cq() → 知道"发完了"
耗时估算
256MB ÷ 50GB/s(400Gbps) ≈ 5.1ms
（对比机内 0.85μs，慢约 6000 倍）
GPU2 收到后做什么
GPU2 把收到的 chunk0（GPU0+GPU1的累加结果）加到自己对应的 chunk0 上：
GPU2.chunk0 = GPU2.chunk0 + 收到的chunk0
= [3,3,3,3] + [3,3,3,3]
= [6,6,6,6]
这就是 Reduce-Scatter 第 2 步的结果。下一步 GPU2 要把 [6,6,6,6] 发给 GPU3（机内，NVLink）。
两种跳对比
这个 chunk 接下来的命运
跳1 GPU0→GPU1 机内 [1,1,1,1] → GPU1加自己的 → [3,3,3,3]
跳2 GPU1→GPU2 跨机 [3,3,3,3] → GPU2加自己的 → [6,6,6,6]
跳3 GPU2→GPU3 机内 [6,6,6,6] → GPU3加自己的 → [10,10,10,10] ✅ 完整
3 跳之后 GPU3 拿到了完整的 chunk0 = [10,10,10,10]（4 个 GPU 的 chunk0 加起来）。 这就是 Reduce-Scatter：N 个 GPU 绕环 N-1 步，每个 GPU 最终拿到 1 块完整的。
接着 AllGather 反向再跑，把完整的块沿环传一圈，所有 GPU 都拿到。
机内跳（跳1、跳3）走 NVLink 白送，跨机跳（跳2）走 IB 是瓶颈——这就是为什么 8 卡 8 NIC 要把跨机段分到 8 张卡上并行走。
11. 一句话收口
ring 是 GPU rank 排成的逻辑环。机内跳走 NVLink 不经过网卡，跨机跳走 IB 网卡经 QP 做 RDMA。一个 chunk 绕环一圈交替经过 NVLink 和 IB，慢的全在 IB 那几跳。8 个 channel 是 8 个拓扑相同的环，跨节点段分别经 8 张 IB 卡中转把带宽打满。ring 被选中是因为它用最少连接数（每GPU只连2邻居）就能把跨机带宽吃满，不是因为它能"拿到数据"——拿到数据是结果，mesh 也能拿到，但 QP 会炸。
