cuda 12.8
sm_90(即 Hopper 架构,如 H100/H200)
5090需要有sm_120架构
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)。有些方案在旧显卡上快,有些方案在新显卡上快。
虽然自动调优是为了性能,但在以下几种情况下它会变成“灾难”:
Illegal Memory Access(非法内存访问)错误。level=0 就是告诉编译器:“别试了,直接给我就近用一个最通用的方案,哪怕它不是最快的。”NCCL配置要注意哪些事项?
NCCL_DEBUG用于调试,NCCL_IB_DISABLE可以禁用InfiniBand支持,如果遇到网络连接问题可以通过这些变量来调整。配置NCCL时的常见问题
NCCL_DEBUG=INFO来获得更多的调试信息在配置和使用NCCL时,有一些重要的环境变量可以帮助优化性能和调试问题:
INFO、WARN、ERROR等不同级别。设置为INFO可以获取详细的调试信息,帮助分析通信过程中的问题。NCCL_IB_DISABLE=1来禁用IB并强制使用Socket通信。NCCL_SOCKET_IFNAME=eth0。NVL(NVLink)、PXB(PCIe跨插槽)、SYS(系统内存)等,以调整通信路径,提升通信效率。1MB。GROUP(组启动)或PARALLEL(并行启动),以优化启动开销,适应不同的集群环境在NCCL的通信中,通常有两种主要的网络通信方式:Socket和InfiniBand(IB)。这两者在通信性能和适用场景上有较大的区别:
日志
在处理NCCL相关问题时,日志是非常重要的工具。通过分析NCCL的日志输出,我们可以获取许多有用的信息来排查问题。例如,通过设置NCCL_DEBUG=INFO,可以看到每个节点的初始化过程、网络通信状况和错误信息。如果训练在某个节点上卡住或者通信效率不佳,可以通过日志查看具体的通信步骤和所使用的通信通道,进而发现是否存在网络配置不一致或者带宽瓶颈等问题。此外,NCCL_DEBUG_SUBSYS=ALL可以进一步细化调试信息,例如打印环路信息和具体的通信操作细节,从而帮助定位问题的根源。
通过这些日志信息,我们可以快速识别出网络不兼容、版本不匹配等问题,从而采取有效的措施进行修复
这表明NCCL在尝试使用InfiniBand(IB)设备时没有找到任何设备。可能是因为系统中没有安装IB设备,或者NCCL没有正确配置以使用这些设备
这些行显示了两个环境变量的配置:
NCCL_IB_DISABLE 设置为0,意味着NCCL被允许使用InfiniBand。NCCL_SOCKET_IFNAME 设置为eth0,指定NCCL应使用名为eth0的网络接口进行通信。由于没有找到IB设备,NCCL转而使用TCP/IP(Socket)网络,并通过接口eth0进行通信。
这些行显示了NCCL为每个GPU设置的CPU亲和性(affinity),这有助于优化GPU和CPU之间的通信。
NVLS(NVIDIA Virtual Link Subsystem)多播支持在该设备上不可用。这可能是因为硬件不支持或者驱动程序配置不正确。
通信器管理(Communicator) :
与 MPI 类似,NCCL 所有通信操作均在 Communicator 上下文环境中执行,参与通信的每个 GPU 均维护一个 Communicator 对象(对应 ncclComm),作为调用 NCCL 的执行载体。用户需首先完成 Communicator 初始化,并明确定义参与通信的 GPU 集合。
ncclComm 中定义了非常全面的信息,包括但不限于:
通信组相关
通信拓扑与连接
算法与性能参数
当所有设备由单一进程/线程管理时,可通过 ncclCommInitAll(位于 http://init.cc)进行 Communicator 创建。在多进程/多线程环境中,各进程需调用 ncclCommInitRank 并共享唯一标识符,以实现跨进程 Communicator 的创建。
通信任务结束后,应正确释放 Communicator 以回收资源。NCCL 提供两个关键函数:
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 操作执行的启动模型,每种方法也有其独特的特性:
通讯协议
NCCL 通过三大硬件组件协调通信过程:
当仅用单个 SM 处理 GPU 工作时,大消息可能导致该 SM 过载,而其他 SM 利用率不足,且无法充分发挥 NVLink 和 IB 等高速链路的带宽。
为避免这一瓶颈,NCCL引入了channel的概念。NCCL把每个集合操作进一步拆分为多个通信channel。每个channel被作为独立的CUDA block 启动,运行在不同的 SM 上;NCCL会把输入缓冲区划分成互不重叠的片段,令各channel并行处理。这种细粒度并行能显著提升总体吞吐量,尤其是针对原本会在单个 SM 上串行的大数据量场景。跨channel分摊工作还能在 NVLink 平台上把流量均匀分布到多块 NIC。
然而,过度使用多 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 间数据传输方式:
为提升带宽利用率,NCCL 采用双二叉树结构——两个树中不存在共用的非叶节点,且至多一个节点在两树中同为叶节点。当节点数为偶数时通过镜像构建第二棵树,奇数时则进行一个位置的偏移。这些拓扑结构在 Communicator 初始化时确立,并在所有集合操作中复用。
对于使用 ncclGroupStart 和 ncclGroupEnd 的分组 P2P 操作,NCCL 会尽可能将每次传输分配到独立 Channel,从而实现多组独立发送与接收操作的并行执行,以此实现传输间的任务级并行
调优模型
NCCL 调优模型对应的实现在 tunning.cc 中,对应 ncclTopoTuneModel 函数,其主要功能和流程如下:
线程数(maxThreads)自动推算:
硬件特性索引与带宽上限选择:
带宽(bandwidth)与延迟(latency)建模:
硬件带宽上限
节点数、每节点 GPU 数
不同协议的效率修正(如 LL128 只在特定架构/拓扑下启用)
特殊算法(如 CollNet、NVLS、PAT)的支持条件和性能修正
网络延迟、PCIe/NVLink/网络等不同链路的延迟模型
算法/协议使能与用户控制:
最终调优参数输出:
为了在集合通信过程中提高数据传输效率,NCCL 采用了多种通信协议。当前主要有三种:Simple、LL(Low Latency)以及LL128。它们各自针对带宽与时延之间的权衡做了不同取舍
Simple 协议
Simple 协议旨在最大化带宽利用率,主要用于大消息传输。它将待传数据切分为较大的块,然后通过多条通信通道并行发送。这种分块策略能够充分发挥网络接口与 GPU 内存系统的高吞吐能力。
为保证内存一致性,Simple 协议依赖内存栅栏(memory fence)来强制数据的顺序与可见性:接收端必须等到一个完整数据块到达后才能访问它。虽然这一做法确保了正确性,但栅栏带来的同步开销显著,尤其在小消息场景中,同步成本往往主导了总传输时间。因此,Simple 协议在大消息时可以接近峰值带宽,而在小消息时则表现出较高的时延。
可能有些人还是不太理解,这里以RDMA为例。由于RDMA Write 操作本身不具备通知能力。为了告知接收方“新数据已准备好”,生产者通常需要采取以下步骤:
LL (Low Latency) Protocol
LL128 Protocol
3.4 协议选择和对比(Protocol Selection and Comparison)
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 干预实现节点间的直接数据移动 。
(1)增大单 QP 的有效数据块大小。
(2)为支持 ECMP 的架构提供路径多样性。
(3)增强整体互连效率。
(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)效应。
集合通信算法是 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 技术实现网络辅助的集合操作:
NVLS 算法旨在利用 NVSwitch 的特性,从而提升集合操作效率。标准 NVLS 与 NVLS Tree 算法均采用 NVLink SHARP 实现节点内 Reduce,但在跨节点处理上存在差异:前者通过 CollNet 及支持 SHARP 的交换机延续 Reduce 过程,后者则采用基于树状结构的扇出传输。
NCCL 通过组合一组底层通信原语来实现高级集合操作。这些原语构成了 NCCL 集合算法的基础,封装了跨 GPU 发送、接收、归约和复制数据等基本操作。常见原语包括:
每个原语代表一种独特的数据移动或计算模式,其命名规范清晰体现了操作顺序。例如 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 中,数据移动的基本单位称为元素,其具体含义取决于集合操作类型:
如上图 Figure 3 展示了该过程的具体实现。图中每个单元格代表 sendBuff 中的一个数据元素。为便于说明,示例设定 channelCount 为 2,chunkCount 为 2,loopCount 为 4。
所有常见的 NCCL 集合通信算法均遵循迭代处理模型,其核心差异在于 GPU 能否对连续循环迭代进行流水线化处理。基于这一特性,算法可分为两类:流水线模式与非流水线模式。
(1)非流水线模式(Non-pipelined Pattern)
特征:每个 GPU 必须完成一次迭代中的所有任务才能开始下一次迭代 。适用于Ring AllReduce、Ring AllGather 和 Ring 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 完成收集 。
nvitop是一个非常全面的NVIDIA-GPU设备运行状况的实时监控工具,它将GPU利用率,显存占比,卡号使用者,CPU利用率,进程使用时间,命令行等等集于一身,并以差异化的颜色进行个性化展示
nvitop是一款交互式NVIDIA-GPU设备性能&资源&进程的实时监测工具。
相比于nvidia-smi命令,nvitop在实时监控GPU设备资源&性能上具备全方位优势: