机器人研究(二十)

CUDA

5090

cuda 12.8

uv run python -c "import torch; print(f'PyTorch Version: {torch.__version__}'); print(f'Supported Architectures: {torch.cuda.get_arch_list()}')"

PyTorch Version: 2.7.1+cu126

Supported Architectures: ['sm_50', 'sm_60', 'sm_70', 'sm_75', 'sm_80', 'sm_86', 'sm_90']

sm_90(即 Hopper 架构,如 H100/H200)

5090需要有sm_120架构

NCCL

NCCL(NVIDIA Collective Communications Library)是NVIDIA发布的一个高效的集体通信库,专为多个GPU之间提供优化的传输效率和简化应用而设计。在现代深度学习中,多机多卡训练已经成为常见的训练模式。然而,由于需要频繁进行大量的数据汇总和分发操作,多机多卡训练面临集体通信的挑战。NCCL为此提供了高性能的通信解决方案,通过接口并行优化的方式,性能远远超过传统的MPI(Message Passing Interface)。

NCCL作为一个专门为GPU集体通信设计的库,主要功能是实现高效的数据计算和数据汇总操作。它能够最大化利用GPU之间的高带宽通信链路,从而实现更高的吞吐量。在多机多卡的训练环境中,NCCL通过高效的通信机制来协调多个GPU之间的数据同步,使得训练过程中的梯度汇聚、参数同步等操作变得更加高效

NCCL主要通过环形拓扑结构来实现高效的集体通信。它支持常见的通信模式,例如AllReduce、Broadcast、Reduce、AllGather等。在多卡训练中,AllReduce是最常用的模式,它用于汇聚所有GPU计算得到的梯度,并将结果分发到每个GPU,以确保模型在每个设备上保持同步。NCCL通过GPU直接通信(GPU Direct)和RDMA(Remote Direct Memory Access)等技术来减少CPU的参与,从而降低延迟和提高数据传输效率

使用NCCL_P2P_DISABLE=1 选项,禁用显卡的P2P通信

NCCL_SHM_DISABLE=1

XLA_FLAGS="--xla_gpu_autotune_level=0"。当你运行 JAX 或 TensorFlow 时,XLA 编译器会发现你的模型里有很多复杂的数学操作(比如卷积或大矩阵乘法)。 对于同一个数学操作,显卡驱动里通常有几十种不同的实现方案(Kernels)。有些方案在旧显卡上快,有些方案在新显卡上快。

  • 默认行为(Level 4):XLA 会在第一次运行(JIT 编译阶段)时,在你的 GPU 上把这几十种方案都“偷偷”跑一遍,测量哪一个最快,然后选定那个最快的方案。
  • 这就是为什么你第一步编译特别慢的原因之一,因为它在做大量的实测

虽然自动调优是为了性能,但在以下几种情况下它会变成“灾难”:

  1. 硬件太新(Blackwell 架构):RTX 5090 的架构非常新。如果 XLA 尝试测试一些针对旧架构优化的方案,可能会触发硬件不支持的操作,直接导致你看到的 Illegal Memory Access(非法内存访问)错误。
  2. 编译器 Bug:在搜索“最优方案”的过程中,编译器可能会误判。设置 level=0 就是告诉编译器:“别试了,直接给我就近用一个最通用的方案,哪怕它不是最快的。”
  3. 显存压力:Autotune 过程本身需要额外申请显存来跑测试代码。如果你的模型已经把显存占满了,Autotune 可能会因为申请不到测试显存而卡死或崩溃

NCCL配置要注意哪些事项?

  1. 拓扑结构的配置:确保多机多卡之间的网络拓扑能够支持高效的通信。NCCL对网络拓扑非常敏感,不同的拓扑结构可能会显著影响通信效率。推荐使用NVLink、PCIe、或高速以太网(如InfiniBand)等高带宽连接。
  2. NCCL环形策略(Ring Strategy):在使用多机多卡时,NCCL采用环形策略来组织GPU之间的通信。需要根据硬件拓扑手动优化环路的设置,以避免潜在的通信瓶颈。
  3. 环境变量:NCCL有一些重要的环境变量可以配置,例如NCCL_DEBUG用于调试,NCCL_IB_DISABLE可以禁用InfiniBand支持,如果遇到网络连接问题可以通过这些变量来调整。
  4. GPU与网络的平衡:确保每个GPU的通信带宽与计算速度匹配。通信瓶颈通常会影响多卡训练的整体性能,因此在配置网络带宽时要尽量与GPU的计算能力相匹配

配置NCCL时的常见问题

  1. 网络不兼容:在多机多卡场景下,不同机器的网络配置不一致会导致NCCL通信失败。例如,某些机器可能没有InfiniBand支持,或者网络接口不一致,可能会导致NCCL初始化失败或通信延迟较高。
  2. CUDA驱动和NCCL版本不匹配:NCCL依赖CUDA驱动,因此必须确保CUDA版本和NCCL版本兼容。不兼容的版本可能会导致通信错误或性能下降。
  3. 环路配置不当:NCCL采用环形拓扑进行通信,如果环路配置不当,会导致某些GPU的通信路径过长,从而影响整体的通信效率。建议根据实际的硬件拓扑手动调整NCCL的环路设置。
  4. 调试困难:NCCL的集体通信问题在调试时往往较为困难,因为涉及多机多卡的网络、驱动、硬件拓扑等多个方面。可以通过设置NCCL_DEBUG=INFO来获得更多的调试信息

在配置和使用NCCL时,有一些重要的环境变量可以帮助优化性能和调试问题:

  1. NCCL_DEBUG:用于调试NCCL的通信问题。可以设置为INFOWARNERROR等不同级别。设置为INFO可以获取详细的调试信息,帮助分析通信过程中的问题。
  2. NCCL_IB_DISABLE:控制是否禁用InfiniBand支持。当系统中存在网络兼容性问题时,可以通过设置NCCL_IB_DISABLE=1来禁用IB并强制使用Socket通信。
  3. NCCL_SOCKET_IFNAME:指定NCCL使用的网络接口名称。在多网络接口的系统中,可以通过设置这个变量来选择特定的接口用于通信,例如NCCL_SOCKET_IFNAME=eth0
  4. NCCL_NET_GDR_LEVEL:控制NCCL是否使用GPU直接通信(GDR,GPUDirect RDMA)。设置较高的级别可以利用GPUDirect RDMA提高通信速度,适合有GPUDirect支持的网络配置。
  5. NCCL_SHM_DISABLE:控制是否禁用共享内存通信。在某些环境下,禁用共享内存可以避免一些因共享内存冲突导致的性能问题。
  6. NCCL_P2P_LEVEL:控制点对点通信的使用级别。可以设置为NVL(NVLink)、PXB(PCIe跨插槽)、SYS(系统内存)等,以调整通信路径,提升通信效率。
  7. NCCL_BUFFSIZE:设置NCCL通信的缓冲区大小。通常设置为更大的值可以提高带宽利用率,但也会占用更多的内存,默认值为1MB
  8. NCCL_LAUNCH_MODE:控制NCCL内核的启动模式。可以设置为GROUP(组启动)或PARALLEL(并行启动),以优化启动开销,适应不同的集群环境

在NCCL的通信中,通常有两种主要的网络通信方式:Socket和InfiniBand(IB)。这两者在通信性能和适用场景上有较大的区别:

  1. 通信性能 Socket:基于以太网的通信方式,通常使用TCP/IP协议栈进行数据传输。Socket通信的优点是网络配置相对简单,适用于广泛的通用网络环境。但是,Socket的带宽和延迟受限于传统以太网的速度,通常较适用于小规模集群或对通信性能要求不太高的场景。 InfiniBand(IB):InfiniBand是一种高速网络技术,具有更高的带宽和更低的延迟。通过使用RDMA(Remote Direct Memory Access)技术,IB能够在不经过CPU的情况下直接在节点间传输数据,从而大大减少了通信延迟。这使得IB成为大规模多机多卡训练的理想选择,尤其是对通信开销敏感的深度学习任务。
  2. 使用场景 Socket:由于其通用性,Socket更适合小型集群、开发测试环境或者预算有限的集群。Socket通信的配置和调试也相对简单,适合初学者或中小规模模型的训练。 InfiniBand(IB):IB主要用于需要大规模并行训练的集群,特别是企业级数据中心和超算中心。它在多机多卡的场景下可以显著减少数据同步的时间,从而提高整体训练的吞吐量。对于大规模模型的训练(例如超过8张GPU的配置),InfiniBand通常是必不可少的。
  3. 配置和兼容性 Socket:基本上所有机器只要有网络接口就可以进行Socket通信,配置简单。但是,其性能瓶颈明显,在进行大规模训练时容易成为系统的瓶颈。 InfiniBand(IB):需要专用的硬件支持(例如Mellanox的网卡和交换机),配置较为复杂且成本较高。此外,还需要特定的驱动和软件栈支持,但在多机环境下,IB的性能优势是无法替代的

日志

在处理NCCL相关问题时,日志是非常重要的工具。通过分析NCCL的日志输出,我们可以获取许多有用的信息来排查问题。例如,通过设置NCCL_DEBUG=INFO,可以看到每个节点的初始化过程、网络通信状况和错误信息。如果训练在某个节点上卡住或者通信效率不佳,可以通过日志查看具体的通信步骤和所使用的通信通道,进而发现是否存在网络配置不一致或者带宽瓶颈等问题。此外,NCCL_DEBUG_SUBSYS=ALL可以进一步细化调试信息,例如打印环路信息和具体的通信操作细节,从而帮助定位问题的根源。

通过这些日志信息,我们可以快速识别出网络不兼容、版本不匹配等问题,从而采取有效的措施进行修复

NCCL INFO NET/IB : No device found.

这表明NCCL在尝试使用InfiniBand(IB)设备时没有找到任何设备。可能是因为系统中没有安装IB设备,或者NCCL没有正确配置以使用这些设备

NCCL INFO NCCL_IB_DISABLE set by environment to 0.
NCCL INFO NCCL_SOCKET_IFNAME set by environment to eth0

这些行显示了两个环境变量的配置:

  • NCCL_IB_DISABLE 设置为0,意味着NCCL被允许使用InfiniBand。
  • NCCL_SOCKET_IFNAME 设置为eth0,指定NCCL应使用名为eth0的网络接口进行通信。
NCCL INFO NET/Socket : Using [0]eth0:10.233.90.231<0>
NCCL INFO Using network Socket

由于没有找到IB设备,NCCL转而使用TCP/IP(Socket)网络,并通过接口eth0进行通信。

NCCL INFO Setting affinity for GPU 2 to 0fffff,ff000000,0fffffff

这些行显示了NCCL为每个GPU设置的CPU亲和性(affinity),这有助于优化GPU和CPU之间的通信。

NCCL INFO NVLS multicast support is not available on dev 2

NVLS(NVIDIA Virtual Link Subsystem)多播支持在该设备上不可用。这可能是因为硬件不支持或者驱动程序配置不正确。

NCCL INFO Using network IBext

代码

通信器管理(Communicator) :

与 MPI 类似,NCCL 所有通信操作均在 Communicator 上下文环境中执行,参与通信的每个 GPU 均维护一个 Communicator 对象(对应 ncclComm),作为调用 NCCL 的执行载体。用户需首先完成 Communicator 初始化,并明确定义参与通信的 GPU 集合。

ncclComm 中定义了非常全面的信息,包括但不限于:

  • 通信组相关

    • rank:当前进程/设备在通信组中的编号。
    • nRanks:通信组总的 rank 数。
    • cudaDev:当前 rank 所在的 CUDA 设备号。
    • localRank/localRanks:当前 rank 在本节点(物理机)内的编号及本节点内的 rank 数,常用于多卡/多机场景。
    • node/nNodes:当前 rank 所在节点编号及总节点数(物理机数)
  • 通信拓扑与连接

    • ncclChannel channels[MAXCHANNELS]:通信 Channel 数组,NCCL 内部用于管理不同算法/协议的通信路径。
    • ncclTopoSystem topo:拓扑结构体指针,描述了当前通信组的硬件拓扑。
    • ncclPeerInfo peerInfo:所有 rank 的硬件和进程信息(如 busId、hostHash 等),用于拓扑和路径选择。
  • 算法与性能参数

    • ncclTopoGraph graphs[NCCL_NUM_ALGORITHMS]:各种集合通信算法(如 Ring、Tree、CollNet、NVLS)的图结构和参数。
    • nChannels:实际用于通信的通道数,影响并发和带宽利用。
    • collNetSupport/nvlsSupport:是否支持 CollNet、NVLS 等高级通信特性

当所有设备由单一进程/线程管理时,可通过 ncclCommInitAll(位于 http://init.cc)进行 Communicator 创建。在多进程/多线程环境中,各进程需调用 ncclCommInitRank 并共享唯一标识符,以实现跨进程 Communicator 的创建。

通信任务结束后,应正确释放 Communicator 以回收资源。NCCL 提供两个关键函数:

  • ncclCommDestroy:安全销毁 Communicator,确保清理前完成所有待处理通信操作。
  • ncclCommAbort:立即终止 Communicator 并取消进行中的操作,适用于错误恢复或意外故障处理场景,可有效避免死锁发生。

2. 集合通信(Collective Communication)

NCCL 提供 5 种集合通信操作:ncclAllReduce、ncclBroadcast、ncclReduce、ncclAllGather 和 ncclReduceScatter。历史原因,NCCL 还包含 ncclBroadcast 的原位操作变体 ncclBcast,以对齐 MPI_Bcast 的行为特性,后续引入了具有独立发送和缓冲区的更通用的 ncclBroadcast,如今 ncclBcast 基本已废弃,仅为兼容 MPI 风格接口而保留。由于兼容性而保留,但是不建议继续使用。

【 NCCL 中并没有 AlltoAll 通信原语,需要通过 ncclSend 和 ncclRecv 实现,其中 ncclSend 和 ncclRecv 是一个 P2P 通信。每个 Rank 都发送 nranks 块数据,同时接收 nranks 块数据就实现了 AlltoAll 的功能】

3.点对点P2P 通信(Point-to-Point Communication):

NCCL 通过 ncclSend 和 ncclRecv 实现点对点通信操作。

4. Group 调用(Group Calls):

为聚合操作并降低系统开销,NCCL 提供 ncclGroupStart 和 ncclGroupEnd 函数,将一系列 NCCL 调用封装为操作组,延迟到组结束时统一执行。Group 操作可包含多个 Send/Recv 调用(如多个 Send/Recv 组合成 SendRecv、All-to-One(Gather)、One-to-All(Scatter) 或 All-to-All;也可是一组集合操作)或一组集合通信操作,组调用会推迟实际执行,直至 ncclGroupEnd,通过该机制可以显著降低启动成本和时延。

启动策略

NCCL 支持 3 种多 GPU 操作执行的启动模型,每种方法也有其独特的特性:

  • 每个 GPU 单个 CPU 进程:提供更精细的进程控制。每个 GPU 绑定到独立进程,关联的 CPU 代码可调度到 Local 的 NUMA 域执行,从而提升数据局部性并降低内存访问时延,对应nccl-test命令示例:
mpirun -np 8 ./build/all_reduce_perf -b 8M -e 128M -f 2 -g 1
#-np 8:启动 8 个进程(对应 8 个 GPU)
#-g 1:每个进程仅管理 1 块 GPU
  • 每个 GPU 对应一个 CPU 线程:当单个进程通过多个线程管理多块 GPU 时,可在进程内部高效共享内存。各 rank 之间(包括 GPU 缓冲区)可以直接访问彼此数据,减少通信过程中的内存拷贝开销。对应nccl-test命令例子如下
./build/all_reduce_perf -b 8M -e 128M -f 2 -g 8 -nthreads 8
#-g 8:进程管理 8 块 GPU
#-nthreads 8:启动 8 个线程(每个线程管理 1 块 GPU)
  • 多 GPU 共享单 CPU 线程:虽然单线程模式会导致内核启动顺序化、并发度下降,但实现简单、CPU 开销极低、执行流程可预测。因而在小规模部署或原型验证场景中,如果易用性比极致性能更重要,这种方式是一个不错的选择。对应nccl-test命令例子如下:
./build/all_reduce_perf -b 8M -e 128M -f 2 -g 8 -nthreads 1
#-g 8:进程管理 8 块 GPU
#-nthreads 1:仅使用 1 个线程控制所有 GPU

通讯协议

NCCL 通过三大硬件组件协调通信过程:

  • GPU:负责执行 Reduce 操作并在缓冲区之间迁移数据。
  • CPU:负责启动计算 Kernel 及 Host 端的协调管理。
  • NIC:承担跨节点数据包传输任务。

当仅用单个 SM 处理 GPU 工作时,大消息可能导致该 SM 过载,而其他 SM 利用率不足,且无法充分发挥 NVLink 和 IB 等高速链路的带宽。

为避免这一瓶颈,NCCL引入了channel的概念。NCCL把每个集合操作进一步拆分为多个通信channel。每个channel被作为独立的CUDA block 启动,运行在不同的 SM 上;NCCL会把输入缓冲区划分成互不重叠的片段,令各channel并行处理。这种细粒度并行能显著提升总体吞吐量,尤其是针对原本会在单个 SM 上串行的大数据量场景。跨channel分摊工作还能在 NVLink 平台上把流量均匀分布到多块 NIC。

  • 每个 Channel 以独立 CUDA Block 形式启动,运行在专属的 SM 上,同时库函数对输入缓冲区进行分区,以确保各 Channel 并行处理并且互不相交,这种细粒度并行机制显著提升了整体吞吐量。
  • 除此之外,还能均衡 NVLink 平台上多 NIC 之间的流量分布(PS:各 Channel 可以使用不同的 NIC 通信)。该策略不仅提高了链路利用率、减少了空闲时间,更能实现 NVLink 、PCIe 及 IB 等互连架构间的负载均衡。

然而,过度使用多 Channel 也可能对网络效率产生负面影响,当单 Channel 数据块小于 NIC 传输层采用的 512KiB FIFO 缓冲区容量时,将会发送未充分填充的缓冲区,进而导致 PCIe 及网络吞吐量下降,特别是启动多 QP(Queue Pair)以实现 ECMP(Equal-Cost Multi-Path Routing)负载均衡的场景。为此,NCCL 采用启发式方法动态减小 Channel 数量(参见 enqueue.cc 中的 calcP2pChunkSize 函数,这里可能是写错了,新版本里对应 calcP2pChannelCount)

NCCL 在 Communicator 初始化阶段会建立一组初始化 Channel,其总数主要由系统拓扑和架构决定。当发起集合通信操作时,NCCL 会动态选择适合对应任务的算法和协议。同时,NCCL 内部调优模型会根据所选策略、当前消息大小、可用带宽及每个 Channel 配置的线程数来确定该操作使用的 Channel 数量。(PS:NCCL 早期版本支持设置 NCCL_NTHREADS 等环境变量来调优通道选择,但新版本都不推荐此类手动调优)

每个 Channel 的逻辑通信拓扑直接影响 GPU 间数据传输方式:

  • Ring 拓扑中:各 GPU 识别其直接前驱和后继节点以形成单向通信环。
  • Tree 拓扑中:各 GPU 记录其父节点和子节点的 Rank 号,构建逻辑通信树。

为提升带宽利用率,NCCL 采用双二叉树结构——两个树中不存在共用的非叶节点,且至多一个节点在两树中同为叶节点。当节点数为偶数时通过镜像构建第二棵树,奇数时则进行一个位置的偏移。这些拓扑结构在 Communicator 初始化时确立,并在所有集合操作中复用。

对于使用 ncclGroupStart 和 ncclGroupEnd 的分组 P2P 操作,NCCL 会尽可能将每次传输分配到独立 Channel,从而实现多组独立发送与接收操作的并行执行,以此实现传输间的任务级并行

调优模型

NCCL 调优模型对应的实现在 tunning.cc 中,对应 ncclTopoTuneModel 函数,其主要功能和流程如下:

线程数(maxThreads)自动推算:

  • 根据算法(如 Ring、Tree、CollNet、NVLS)和协议(Simple、LL、LL128),结合带宽、通道数、PCIe带宽等,自动为每种算法/协议选择合适的线程数。
  • 支持通过环境变量(如 NCCL_NTHREADS、NCCL_LL128_NTHREADS)进行覆盖。

硬件特性索引与带宽上限选择:

  • 根据 GPU 架构(Volta/Ampere/Hopper/Blackwell)和节点数,选择合适的带宽上限表(如下图所示的 llMaxBws、perChMaxTreeBws 等)。
  • 单节点时主要看 GPU 类型,多节点时还考虑 CPU 厂商(Intel/AMD)。

带宽(bandwidth)与延迟(latency)建模:

  • 对每种集合通信操作(如 AllReduce、Broadcast、AllGather、ReduceScatter)、每种算法、每种协议,计算理论带宽和延迟。

硬件带宽上限

节点数、每节点 GPU 数

不同协议的效率修正(如 LL128 只在特定架构/拓扑下启用)

特殊算法(如 CollNet、NVLS、PAT)的支持条件和性能修正

网络延迟、PCIe/NVLink/网络等不同链路的延迟模型

算法/协议使能与用户控制:

  • 支持通过环境变量 NCCL_ALGO、NCCL_PROTO 精细控制哪些算法/协议可用。
  • 自动禁用不支持的算法(如单节点禁用 CollNet/NVLS Tree,多节点无 NVSwitch 禁用 CollNet Direct)。
  • LL128 协议只在特定硬件和配置下默认启用,防止数据错误。

最终调优参数输出:

  • 计算并输出每种算法/协议/操作的最终带宽、延迟、线程阈值等参数,供 NCCL 内部调度和选择最优通信方案。
  • 支持通过环境变量 NCCL_THREAD_THRESHOLDS 覆盖线程阈值

通信协议

为了在集合通信过程中提高数据传输效率,NCCL 采用了多种通信协议。当前主要有三种:Simple、LL(Low Latency)以及LL128。它们各自针对带宽与时延之间的权衡做了不同取舍

Simple 协议

Simple 协议旨在最大化带宽利用率,主要用于大消息传输。它将待传数据切分为较大的块,然后通过多条通信通道并行发送。这种分块策略能够充分发挥网络接口与 GPU 内存系统的高吞吐能力。

为保证内存一致性,Simple 协议依赖内存栅栏(memory fence)来强制数据的顺序与可见性:接收端必须等到一个完整数据块到达后才能访问它。虽然这一做法确保了正确性,但栅栏带来的同步开销显著,尤其在小消息场景中,同步成本往往主导了总传输时间。因此,Simple 协议在大消息时可以接近峰值带宽,而在小消息时则表现出较高的时延。

可能有些人还是不太理解,这里以RDMA为例。由于RDMA Write 操作本身不具备通知能力。为了告知接收方“新数据已准备好”,生产者通常需要采取以下步骤:

  1. 写入数据: 使用一个或多个 RDMA Write 操作将实际数据写入接收方的内存缓冲区。
  2. 发出内存屏障(Fence): 发送方发出一个带有 IBV_SEND_FENCE 标志的 WR (通常是另一个 RDMA Write 或一个 No-Op 操作)。这个屏障确保之前所有的 RDMA Write (步骤 1 的数据写入) 必须在之后的任何 RDMA Write 之前完成,并且其效果对接收方可见。
  3. 更新标志位(Doorbell): 发送方执行 另一个 RDMA Write 操作,去修改接收方内存中的一个特定位置(标志位,如一个 ready 变量)。这个标志位用来通知接收方:“数据已经准备好了,你可以读了”

LL (Low Latency) Protocol

  • 为低延迟优化:为解决 Simple 协议的延迟问题,NCCL 引入了 LL 协议,该协议专为小消息优化,因为小消息场景下带宽通常未被充分利用 。
  • 基于标志的轻量级同步:LL 协议不依赖内存屏障,而是使用轻量级的基于标志的同步(flag-based synchronization)。一个小的标志(flag)与数据一同传输以表示其有效性,使接收方一旦数据可用就能立即处理,无需昂贵的内存屏障 。
  • 实现细节与性能权衡:每次 LL 协议传输包含 4 字节数据和 4 字节标志,通过 8 字节原子操作一同发送 。这种方法显著降低了同步开销 。然而,它强制要求中间缓冲区必须位于主机内存中,以便 CPU 可以轮询标志并检测数据何时准备好通过 NIC 发送 。这是因为通过 PCIe 轮询 GPU 内存比访问 DRAM 慢得多,并且需要显式同步来确保数据在主机上的可见性 。此设计虽然实现了低延迟,但也
  • 阻止了 GPUDirect RDMA 的使用,严重限制了带宽 。因此,LL 协议通常只能达到峰值带宽的 25-50% ,仅在延迟至关重要且带宽利用率次要的小数据量传输中被优先选择

LL128 Protocol

  • 兼顾低延迟与高带宽:LL128 协议在保持 LL 协议低延迟特性的同时,显著提高了带宽效率,尤其是在 NVLink 等高性能互连上 。它同样采用基于标志的同步来避免内存屏障,但以 128 字节为单位传输数据 。
  • 传输单元与性能:在这 128 字节中,120 字节用于数据,8 字节保留给标志,使得该协议能够利用大约 95% 的峰值带宽 。在网络路径上,LL128 类似于 Simple 协议,发送方 GPU 会聚合一个相对较大的数据块,然后通知 CPU 发送 。虽然这种基于块的聚合限制了跨节点的流水线操作,但由于其较小的传输粒度,LL128 仍然能在节点内部实现细粒度的流水线 。
  • 硬件依赖性:LL128 对硬件有更严格的要求,它依赖于 128 字节的原子写操作,这些操作不能被内存系统或互连拆分或重排 。在由于 PCIe 限制或其他架构约束而无法保证此类操作的系统中,NCCL 会禁用 LL128 以避免数据损坏 。

3.4 协议选择和对比(Protocol Selection and Comparison)

  • 动态选择机制:NCCL 在运行时根据用户设置(NCCL_PROTO)、集合算法以及内部性能启发式动态地在 Simple、LL 和 LL128 协议中进行选择 。若未明确指定,NCCL 会使用一个调优模型,该模型综合考虑系统拓扑、GPU 架构、消息大小和预定义的性能指标来选择最佳的算法-协议对 。该选择受到内存等资源的限制 。通常,LL/LL128 被用于小消息以降低延迟,而 Simple 被用于大消息以最大化吞吐量

数据传输策略和传输层

NCCL 根据通信是在单个节点内部(intra-node)还是跨多个节点(inter-node)发生,采用不同的数据传输策略和传输机制 。每种传输机制均针对特定硬件架构与互连类型进行优化,以此支撑可扩展的集合通信操作

NCCL 采用层次化架构来实现节点内通信,优先选择同一物理机内 GPU 间延迟最低、带宽最高的传输路径。该策略深度依赖 NVIDIA 的 GPUDirect P2P 技术,使得 GPU 能够直接访问彼此显存,无需经由 CPU 系统内存中转

P2P Transport (P2P传输):P2P 传输是核心策略,主要在 src/transport/http://p2p.cc 中管理 。当 GPU 通过 NVIDIA NVLink 互连时,NCCL 优先使用此路径实现 GPUDirect P2P 。若 NVLink 不可用,NCCL 会回退到通过 PCIe 总线使用 GPUDirect P2P 。

  • P2P_DIRECT 优化:当通信的 rank 属于同一进程时,会启用 P2P_DIRECT 模式 。此模式通过两种方式显著提高效率:首先,它通过在同一地址空间内使用直接的 GPU 内存指针,绕过了 IPC 句柄的需求 ;其次,它通过使用direct Send 和 direct Recv 等原语,直接在源和目标缓冲区之间传输数据,而不是通过中间 FIFO 缓冲区,从而消除了一个中间数据拷贝 。

  • Shared Memory (SHM) Transport (共享内存传输):当直接 P2P 通信不可用或性能不佳时(例如,跨 CPU 插槽的 PCIe P2P 性能下降),NCCL 可能会利用 SHM 传输 。SHM 通过系统内存路由流量,利用 CPU 优化的 PCIe-内存和内存-PCIe 传输来避免性能问题 。

  • 通过NIC进行节点内通信:在某些多插槽系统中,如果每个 GPU 位于独立的 CPU 插槽上,且各自拥有支持 GPUDirect RDMA 的本地 NIC,NCCL 可能会使用 NIC 进行节点内通信 。这种 GPU-NIC-NIC-GPU 路径利用 PCIe 带宽,避免了 CPU 互连的瓶颈 。此行为由 NCCL 的拓扑感知逻辑决定,并可通过NCCL_CROSS_NIC 等环境变量控制

  • 1) Socket-Based Communication (基于套接字的通信):当网络接口不支持 RDMA 时,NCCL 使用在 transport/http://net_socket.cc 中实现的套接字传输 。在这种模式下,数据需从 GPU 拷贝到主机固定内存(pinned memory)中,然后通过标准套接字调用发送 。这会产生额外的 PCIe 总线拷贝开销 。

  • 2) IB Verbs Transport (IB Verbs 传输):对于 InfiniBand 或 RoCE 等高性能网络,NCCL 使用在 http://net_ib.cc 中实现的 IB 传输 。它利用 RDMA 功能以最少的 CPU 干预实现节点间的直接数据移动 。

    • a) The GPUDirect RDMA (GDRDMA) Optimization:GDRDMA 是 IB 传输中的一个关键优化,它使 NIC 能够直接访问 GPU 内存,从而消除了主机内存中转 。这仅在 NIC 和 GPU 连接到同一个 PCIe 交换机时使用 。CPU 代理线程使用如nv_peer_mem (【Developing a Linux Kernel Module using GPUDirect RDMA, 2025, NVIDIA Corporation】) 或 Linux DMA-BUF 子系统 (【Buffer Sharing and Synchronization (dma-buf), 2025, http://kernel.org】) 等机制将 GPU 内存注册到 NIC,使其可以直接进行 RDMA 读写 。
    • b) Per-peer Multi-channel Connections (每对等体多信道连接):为提高带宽利用率,IB 传输默认情况下为每个远程 GPU 和每个 NIC 实例化 2 个逻辑信道(由 NCHANNELS_PER_NET_PEER 参数控制),每条逻辑 Channel 维护独立的 ncclIbSendComm 结构体,内含一组完整的 IB QP 集合。运行时,Host 侧网络代理在发起 ncclNet->isend() 调用时交替使用两个 sendComm 句柄,从而将流量分流至不同 QP 集合。这种轻量级轮询策略通过三重机制提升效能,且不会引入额外的 GPU 端状态维护开销:

(1)增大单 QP 的有效数据块大小。

(2)为支持 ECMP 的架构提供路径多样性。

(3)增强整体互连效率。

    • c) QP Layout (QP布局):对于每对 rank,RDMA 插件会建立两个可靠连接(RC)QP,每个方向一个 ,实现双向通道:

(1) forward QP 负责大数据流传输:代理端发起一个或多个 RDMA_WRITE 工作请求,将用户数据直接推送至对端缓冲区,最终以零字节 RDMA_WRITE_WITH_IMM 操作收尾。该请求的 immediate 数据字段编码了传输总量,接收方通过轮询此字段确认传输完成。

(2) reverse QP 仅传输小的 clear-to-send(CTS)消息,通过单次 RDMA_WRITE 操作通知远程缓冲区地址、rkeys 及标签信息。

(3) 虽然理论上可通过单 QP 实现功能复用,但将 CTS 隔离至独立 Channel 能有效分离延迟敏感的控制流量与带宽密集型数据流,从而最大限度降低网络传输中的队首阻塞(head-of-line blocking)效应。

    • d) Local Flush with Loop-back RDMA READ (使用环回RDMA READ进行本地刷新):启用 GDRDMA 时,发送方必须确保所有未完成的 PCIe 写操作都已到达 GPU 内存 。NCCL 通过在最后一次接收完成后发出一个虚拟的RDMA_READ 来实现这一点 。这个RDMA_READ 操作被发送到一个连接自身的“刷新”QP上,因此读操作从未离开主机,但 Verbs 层会等待先前的 PCIe 写操作完成,从而提供了一个廉价的排序屏障

NCCL 集合通信算法

集合通信算法是 NCCL 的核心,它实现了 GPU 间高效、同步的通信。NCCL 通过将每个集合操作分解为底层通信原语,并将其分配到多个并行 Channel 中来实现这些算法。算法选择取决于具体的集合操作及相关执行参数,如消息大小和拓扑结构。

NCCL 提供 6 种算法,但并非每种算法均适合每种协议。如下图 Table III 所示,对应 NCCL 2.19 中 5 种集合通信操作支持的算法与通信协议,该数据源自 src/device 目录下的相应头文件。

  • NVLS 与 CollNet 是专为优化 AllReduce 性能设计的特殊算法,其中 NVLS 还通过利用特定硬件能力支持 ReduceScatter 和 AllGather 操作。

  • CollNet 算法适用于网络基础设施可直接参与集合运算的场景,例如采用 NVIDIA SHARP 技术时,可将 Reduce 运算或其他部分集合计算卸载至网络交换机执行,从而减少数据迁移与延迟。

  • CollNet 算法依托 NVIDIA SHARP 技术实现网络辅助的集合操作:

    • CollNet Direct 支持节点内 All2All 通信。
    • CollNet Chain 则采用线性拓扑排列 GPU,沿链式结构进行上行 Reduce 与下行 Broadcast。
  • NVLS 算法旨在利用 NVSwitch 的特性,从而提升集合操作效率。标准 NVLS 与 NVLS Tree 算法均采用 NVLink SHARP 实现节点内 Reduce,但在跨节点处理上存在差异:前者通过 CollNet 及支持 SHARP 的交换机延续 Reduce 过程,后者则采用基于树状结构的扇出传输。

NCCL 通过组合一组底层通信原语来实现高级集合操作。这些原语构成了 NCCL 集合算法的基础,封装了跨 GPU 发送、接收、归约和复制数据等基本操作。常见原语包括:

  • send
  • recv
  • recvReduceSend
  • recvCopySend
  • recvReduceCopySend
  • 同时还包括相应的 "direct" 变体。

每个原语代表一种独特的数据移动或计算模式,其命名规范清晰体现了操作顺序。例如 recvReduceSend 表示 GPU 从对端接收数据,与本地缓冲区执行归约操作,并将结果发送至下一 GPU 的步骤。在执行过程中,NCCL 运行时通过循环步骤迭代调度这些原语,从而实现对不同算法、拓扑结构和传输层的灵活协调。

NCCL 原语的具体行为还受所选通信协议影响。根据采用 Simple、LL 或 LL128 协议的不同,同步机制、缓冲区管理和传输粒度会存在差异。需要特别指出的是,这些底层原语针对源节点和目标节点数量固定且较少的集合操作(如 Ring 和 Tree 拓扑,通常涉及单一源节点和目标节点,某些 Tree 拓扑最多三个节点)进行了深度优化。虽然这种方法使许多标准集合算法能实现高效运行,但对于 All2All 等需要处理 N 个源节点和 N 个目标节点的通信模式则效率欠佳

NCCL 处理集合操作时,首先将用户输入数据分配到可用的通信 Channel 中,从而实现 Channel 级并行。每个通道负责输入数据的一个连续片段,其范围由元素总数(count)与 Channel 数量共同决定。如下图 Figure 3 所示,数据被划分为若干区域,各 Channel(如 Channel 0 和 Channel 1)独立处理其分配区段。各 Channel 的起始索引由 workOffset 确定,处理规模由 channelCount 指定

为优化数据传输与计算效率,NCCL 为每个 Channel 分配固定大小的缓冲区,其容量取决于所选通信协议(Simple、LL 或 LL128,详见下图 Table IV)。若 Channel 的数据区域超出缓冲区容量,NCCL 将数据分解为若干外层循环迭代。每次迭代处理不超过缓冲区大小的数据段(每次迭代处理 loopCount 个元素),Channel 通过多次循环完成全部分配元素的处理

在每次外层循环迭代中,NCCL 采用流水线技术:将 Channel 缓冲区划分为固定数量的 Slot(通常为 8 个,由 NCCL_STEPS 参数设定)。每个 Slot 可独立推进通信与计算的不同阶段,实现数据传输与规约/复制操作的流水线重叠。每个基础步骤处理一个数据块(含 chunkCount 个元素,循环末块为 lastChunkCount),并将其映射至缓冲区 Slot。这种分块机制确保通信 Channel 持续饱和,通过新数据块与进行中操作的重叠实现最大吞吐量。

在 NCCL 中,数据移动的基本单位称为元素,其具体含义取决于集合操作类型:

  • 对于 ncclAllGather 和 ncclBroadcast 操作,每个元素即单个字节,因为这些操作的核心目标是高效移动与拼接数据。字节级粒度赋予 NCCL 在数据打包与传输上的灵活性,使其独立于底层数据类型。
  • 对于 ncclAllReduce、ncclReduceScatter 和 ncclReduce 操作,每个元素对应着用户自定义数据类型(如 float 或 int),因为这些运算需要基于数据类型层面才有意义的算术规约操作。

如上图 Figure 3 展示了该过程的具体实现。图中每个单元格代表 sendBuff 中的一个数据元素。为便于说明,示例设定 channelCount 为 2,chunkCount 为 2,loopCount 为 4。

  • Channel 0 从其 workOffset 起始,以 loopCount 为循环迭代步长处理元素,并进一步将其分解为 chunkCount 大小的数据块。
  • Channel 1 在其对应区域遵循相同处理逻辑

所有常见的 NCCL 集合通信算法均遵循迭代处理模型,其核心差异在于 GPU 能否对连续循环迭代进行流水线化处理。基于这一特性,算法可分为两类:流水线模式与非流水线模式。

(1)非流水线模式(Non-pipelined Pattern)

  • 特征:每个 GPU 必须完成一次迭代中的所有任务才能开始下一次迭代 。适用于Ring AllReduceRing AllGatherRing ReduceScatter

  • a) Ring AllReduce:该算法在一个循环中包含 2k−1 个步骤,结合了 ReduceScatter 和 AllGather 两个阶段(k为GPU数量) 。

    • Reduce-Scatter 阶段 (步骤 0 到 k-1):每个 GPU 最初发送自己的一段数据(send),然后在接下来的 k−2 步中重复执行 recvReduceSend(接收、归约、发送),最后一步执行 recvReduceCopySend(接收、完成归约、复制到输出缓冲区并发送)。

    • All-Gather 阶段 (步骤 k 到 2k-2):在此阶段,每个 GPU 接收完整的、已归约的数据块。GPU 重复执行 recvCopySend(接收、复制到输出、转发),最后一步执行 recv 完成收集 。

Nvtop

nvitop是一个非常全面的NVIDIA-GPU设备运行状况的实时监控工具,它将GPU利用率,显存占比,卡号使用者,CPU利用率,进程使用时间,命令行等等集于一身,并以差异化的颜色进行个性化展示

nvitop是一款交互式NVIDIA-GPU设备性能&资源&进程的实时监测工具。

相比于nvidia-smi命令,nvitop在实时监控GPU设备资源&性能上具备全方位优势:

  • 以更美观的颜色,和更直观的进度条实时展示某块GPU卡所处进程的GPU&CPU内存以及利用率占比
  • 作为资源监控器,它包括如下功能:树状视图、环境变量查看、进程过滤、进程指标检测等
  • 可追踪某个单项进程在GPU&CPU上内存和利用率占比的历史纪录,并利用Bar直观展示
  • 可直观展示某块GPU的使用者、使用时间、使用命令行、GPU和CPU占用率记录
  • 可使用工具提供的API搭建自定义监控工具

https://nvitop.readthedocs.io/en/latest/