下面的示意图以 2 个 node、4 个 CUDA device、4 个 rank 为例，展示 NCCL 中各对象的归属关系、通信任务的提交过程，以及运行时的数据传输机制。每个 node 上有一个 Process，分别管理两个 CUDA device；四个 rank 共同组成一个 Communicator，通过多个 Channel 并行处理通信数据。整体分为六个层次，同一列尽量对应同一个 rank 及其资源，将逻辑身份、软件对象与物理设备联系起来。
1. 身份与成员层：确定哪些成员参与通信。
Communicator 定义一组通信参与者及其通信上下文。rank 是成员在该 Communicator 中的逻辑编号，范围为 0 … nRanks−1。每个 rank 通过本地的 ncclComm_t 对象保存自身状态，并关联一个 CUDA device。图中的 comm0～comm3 是同一个 Communicator 在四个 rank 上的本地对象。
UniqueId 是初始化 Communicator 时使用的引导标识，其数据类型为 ncclUniqueId。在图示的 ncclCommInitRank() 初始化方式下，由一个参与者调用 ncclGetUniqueId() 生成 UniqueId，再通过 MPI、Socket 等外部机制分发给所有参与者。因此，图中的四个本地对象使用相同的 UniqueId；独立初始化另一个 Communicator 时，应重新生成 UniqueId。
UniqueId 与 Communicator 并非在所有创建方式下都一一对应：ncclCommInitRankScalable() 可以使用多个 UniqueId 初始化一个 Communicator；ncclCommSplit() 则可以从已有 Communicator 派生新 Communicator，无须用户重新生成 UniqueId。
2. Host 调度层：组织 API 调用并生成执行计划。
Group 是由 ncclGroupStart() 和 ncclGroupEnd() 包围、属于当前 Host thread 的调用范围，可以组织一个线程对多个 CUDA device 的操作，也可以聚合通信调用。最外层 ncclGroupEnd() 统一推进所收集的操作。图中分别展示初始化 Group 和通信 Group，两者属于不同的调用范围。
图中的 job 特指 NCCL 内部的 Host 任务，例如初始化使用的 ncclAsyncJob。它包含执行函数与任务状态，可由当前线程或 worker 执行。
planner 负责将 API 提交的通信任务编排为执行计划。对于图示的常规 kernel 路径，它选择 algorithm / protocol，分配 Channel / chunk，生成 ncclKernelPlan 和 work descriptors，并准备所需的 Proxy 操作。随后由 launch 路径将选定的 kernel 和参数提交给 GPU。这里的“转换”主要是选择和编排已有 kernel，并不意味着将每次 API 调用编译成一个新 kernel；多个调用也可能被组织进同一个执行计划。
3. 本地资源层：展示每个 rank 拥有哪些对象。
每个本地 Communicator 都维护自己的 planner、work 存储和 Channel 等资源。图中采用 workFifo 路径：同一个本地 Communicator 的多个 Channel 共享这份 work 存储，GPU kernel 从中读取执行描述。
每个 Channel 保存自己的 Ring、Tree 等拓扑信息，以及面向通信伙伴的 peers[p] 状态。其中，send / recv 表示发送方向和接收方向的连接资源，包含 buffer、进度信息及 transport 所需的连接状态。四列结构相同，说明各 rank 都维护相应的本地资源。
4. 并行通路层：展示通信数据如何分配给 Channel。
Channel 是 Communicator 内部组织并行通信的逻辑通路。一次通信操作的数据可以划分为多个 chunk，分配给不同 Channel 处理。图中的两个 Channel 使用不同的 Ring 顺序，并分别突出其负责的数据块。
Ring 表示 rank 之间的环形通信关系，Tree 表示树形通信关系。相关拓扑在初始化过程中建立，planner 根据具体操作选择 Channel 并分配数据。Channel 是软件中的逻辑通路，与物理 GPU 互联或 RNIC 之间没有固定的一一对应关系。
5. 队列与运行时层：协调执行描述、数据传输与进度。
workFifo 是 Host planner 向 GPU kernel 提供 work descriptors 的环形存储区，记录 buffer 地址、数据量和 Channel 范围等执行信息。图中的 ncclDevWorkColl 展示了一种集合通信 work descriptor。workFifoProduced 和 workFifoConsumed 跟踪空间的使用与回收，图示路径通过 CUDA event 和 callback 确认相应工作完成。FIFO、Args、Persistent 是不同的 work 存储方式。
connFifo 是连接运行时使用的 slot 元数据数组，GPU kernel 与 CPU Proxy 通过它协作。Proxy 是在 Host 上推进传输工作的执行组件。connFifo 的生产者与消费者需要按具体字段判断：在图示的 NET + Simple 路径中，Send 的 size 由 GPU kernel 写入、CPU Proxy 读取；Shared buffer 模式的 offset 由 CPU Proxy 写入、GPU kernel 读取。因此，connFifo 中包含双向交换的信息。真正的 payload 存放在独立的 buffs[protocol] 等数据 buffer 中。
step 表示推进步骤，通过 step % NCCL_STEPS 选择环形 slot；head / tail 是连接中独立的进度字段，用于协调数据就绪与空间复用，不属于 connFifo slot。Send 路径中，Proxy 需要结合 tail 判断数据是否准备完成，不能只看 size。Recv 路径中，RNIC 写入 payload，Proxy 确认接收完成并完成所需的 flush 后更新 tail，GPU kernel 再消费数据。具体机制随 transport、protocol 和 buffer 模式而变化。
6. 物理映射层：对应实际设备与互联。
图中的 node 表示一台物理主机，CUDA device 表示进程可见的 CUDA 设备。设备编号属于进程的可见设备空间，因此两个 node 上都可以出现 CUDA device 0；而 rank 编号属于 Communicator，两者含义不同。
GPU 内部的 SM 执行计算与通信 kernel，CE 表示 Copy Engine，HBM 存放设备数据。节点内的 GPU 通过 NVLink 互联，并通过 PCIe 与 RNIC 交换数据。RNIC 是支持 RDMA 的网络接口设备，节点之间通过图中所示的 IB / RoCE Network 通信。
)
