一、背景

筆者之前的文章中詳細(xì)介紹過 NCCL 初始化階段的拓?fù)浣?、通信路徑?jì)算和優(yōu)化等工作，也介紹過一些 LLM 訓(xùn)練和推理中對(duì) NCCL 的優(yōu)化工作。本文中，借著一篇新的論文具體介紹一下 NCCL 的內(nèi)部設(shè)計(jì)原理和運(yùn)行機(jī)制。

對(duì)應(yīng)的論文：[2507.04786] Demystifying NCCL: An In-depth Analysis of GPU Communication Protocols and Algorithms [1]NCCL 對(duì)應(yīng)的代碼庫：GitHub - NVIDIA/nccl: Optimized primitives for collective multi-GPU communication [2]

二、摘要

NCCL 是實(shí)現(xiàn)大規(guī)模 GPU 集群高性能集合通信操作的關(guān)鍵軟件層，其核心特性在于低時(shí)延和高帶寬。盡管該庫已開源并提供 API 文檔，但其內(nèi)部設(shè)計(jì)原理仍存在顯著的不透明性。通信 Channel 的協(xié)調(diào)機(jī)制、協(xié)議選擇策略以及跨設(shè)備/跨節(jié)點(diǎn)內(nèi)存移動(dòng)的處理方式尚未得到充分解析，導(dǎo)致性能分析與瓶頸定位困難。

本文中，作者對(duì) NCCL 展開系統(tǒng)性研究，重點(diǎn)剖析了其通信協(xié)議(Simple/LL/LL128)、節(jié)點(diǎn)內(nèi)與節(jié)點(diǎn)間的傳輸控制機(jī)制、以及基于 Ring 和 Tree 的通信算法。

本文的研究也構(gòu)成了作者另一篇文章 [2505.08936] ATLAHS: An Application-centric Network Simulator Toolchain for AI, HPC, and Distributed Storage [3] 的理論基礎(chǔ)。

三、引言

3.1 NCCL API

NCCL 面向用戶主要提供 4 大模塊：

Communicator 管理：與 MPI 類似，NCCL 所有通信操作均在 Communicator 上下文環(huán)境中執(zhí)行，參與通信的每個(gè) GPU 均維護(hù)一個(gè) Communicator 對(duì)象(對(duì)應(yīng) ncclComm)，作為調(diào)用 NCCL 的執(zhí)行載體。用戶需首先完成 Communicator 初始化，并明確定義參與通信的 GPU 集合。

ncclComm 中定義了非常全面的信息，包括但不限于：

1.通信組相關(guān)

• rank：當(dāng)前進(jìn)程/設(shè)備在通信組中的編號(hào)。
• nRanks：通信組總的 rank 數(shù)。
• cudaDev：當(dāng)前 rank 所在的 CUDA 設(shè)備號(hào)。
• localRank/localRanks：當(dāng)前 rank 在本節(jié)點(diǎn)(物理機(jī))內(nèi)的編號(hào)及本節(jié)點(diǎn)內(nèi)的 rank 數(shù)，常用于多卡/多機(jī)場(chǎng)景。
• node/nNodes：當(dāng)前 rank 所在節(jié)點(diǎn)編號(hào)及總節(jié)點(diǎn)數(shù)(物理機(jī)數(shù))。

2.通信拓?fù)渑c連接

• ncclChannel channels[MAXCHANNELS]：通信 Channel 數(shù)組，NCCL 內(nèi)部用于管理不同算法/協(xié)議的通信路徑。
• ncclTopoSystem topo：拓?fù)浣Y(jié)構(gòu)體指針，描述了當(dāng)前通信組的硬件拓?fù)洹?/li>
• cclPeerInfo peerInfo：所有 rank 的硬件和進(jìn)程信息(如 busId、hostHash 等)，用于拓?fù)浜吐窂竭x擇。

3.算法與性能參數(shù)

• ncclTopoGraph graphs[NCCL_NUM_ALGORITHMS]：各種集合通信算法(如 Ring、Tree、CollNet、NVLS)的圖結(jié)構(gòu)和參數(shù)。
• nChannels：實(shí)際用于通信的通道數(shù)，影響并發(fā)和帶寬利用。
• collNetSupport/nvlsSupport：是否支持 CollNet、NVLS 等高級(jí)通信特性。

當(dāng)所有設(shè)備由單一進(jìn)程/線程管理時(shí)，可通過 ncclCommInitAll(位于 init.cc)進(jìn)行 Communicator 創(chuàng)建。在多進(jìn)程/多線程環(huán)境中，各進(jìn)程需調(diào)用 ncclCommInitRank 并共享唯一標(biāo)識(shí)符，以實(shí)現(xiàn)跨進(jìn)程 Communicator 的創(chuàng)建。

通信任務(wù)結(jié)束后，應(yīng)正確釋放 Communicator 以回收資源。NCCL 提供兩個(gè)關(guān)鍵函數(shù)：

• ncclCommDestroy：安全銷毀 Communicator，確保清理前完成所有待處理通信操作。
• ncclCommAbort：立即終止 Communicator 并取消進(jìn)行中的操作，適用于錯(cuò)誤恢復(fù)或意外故障處理場(chǎng)景，可有效避免死鎖發(fā)生。

集合通信：NCCL 提供 5 種集合通信操作：ncclAllReduce、ncclBroadcast、ncclReduce、ncclAllGather 和 ncclReduceScatter。歷史原因，NCCL 還包含 ncclBroadcast 的原位操作變體 ncclBcast，以對(duì)齊 MPI_Bcast 的行為特性，后續(xù)引入了具有獨(dú)立發(fā)送和緩沖區(qū)的更通用的 ncclBroadcast，由于兼容性而保留，但是不建議繼續(xù)使用。

P2P 通信：NCCL 通過 ncclSend 和 ncclRecv 實(shí)現(xiàn)點(diǎn)對(duì)點(diǎn)通信操作。

Group 調(diào)用：為聚合操作并降低系統(tǒng)開銷，NCCL 提供 ncclGroupStart 和 ncclGroupEnd 函數(shù)，將一系列 NCCL 調(diào)用封裝為操作組，延遲到組結(jié)束時(shí)統(tǒng)一執(zhí)行。Group 操作可包含多個(gè) Send/Recv 調(diào)用(用于模擬 SendRecv、All2One、One2All 和 All2All 通信操作)或一組集合通信操作，該機(jī)制能顯著降低啟動(dòng)開銷與延遲。

3.2 啟動(dòng)策略

NCCL 支持 3 種多 GPU 操作執(zhí)行的啟動(dòng)模型，每種方法也有其獨(dú)特的特性：

• 每個(gè) GPU 單個(gè) CPU 進(jìn)程：提供更精細(xì)的進(jìn)程控制。每個(gè) GPU 綁定到獨(dú)立進(jìn)程，關(guān)聯(lián)的 CPU 代碼可調(diào)度到 Local 的 NUMA 域執(zhí)行，從而提升數(shù)據(jù)局部性并降低內(nèi)存訪問時(shí)延。
• 每個(gè) GPU 單個(gè) CPU 線程：可實(shí)現(xiàn)高效的進(jìn)程內(nèi)共享內(nèi)存訪問，減少通信過程中的內(nèi)存拷貝開銷。
• 多 GPU 共享單 CPU 線程：實(shí)現(xiàn)簡(jiǎn)單、CPU 開銷小，并且具備確定性的執(zhí)行順序，適合小規(guī)模場(chǎng)景或原型環(huán)境。

3.3 通信 Channel

NCCL 通過三大硬件組件協(xié)調(diào)通信過程：

• GPU：負(fù)責(zé)執(zhí)行 Reduce 操作并在緩沖區(qū)之間遷移數(shù)據(jù)。
• CPU：負(fù)責(zé)啟動(dòng)計(jì)算 Kernel 及 Host 端的協(xié)調(diào)管理。
• NIC：承擔(dān)跨節(jié)點(diǎn)數(shù)據(jù)包傳輸任務(wù)。

當(dāng)僅用單個(gè) SM 處理 GPU 工作時(shí)，大消息可能導(dǎo)致該 SM 過載，而其他 SM 利用率不足，且無法充分發(fā)揮 NVLink 和 IB 等高速鏈路的帶寬。

為了消除這一瓶頸，NCCL 將每個(gè)集合通信操作細(xì)分為多個(gè)通信 Channel。

• 每個(gè) Channel 以獨(dú)立 CUDA Block 形式啟動(dòng)，運(yùn)行在專屬的 SM 上，同時(shí)庫函數(shù)對(duì)輸入緩沖區(qū)進(jìn)行分區(qū)，以確保各 Channel 并行處理并且互不相交，這種細(xì)粒度并行機(jī)制顯著提升了整體吞吐量。
• 除此之外，還能均衡 NVLink 平臺(tái)上多 NIC 之間的流量分布(PS：各 Channel 可以使用不同的 NIC 通信)。該策略不僅提高了鏈路利用率、減少了空閑時(shí)間，更能實(shí)現(xiàn) NVLink 、PCIe 及 IB 等互連架構(gòu)間的負(fù)載均衡。

然而，過度使用多 Channel 也可能對(duì)網(wǎng)絡(luò)效率產(chǎn)生負(fù)面影響，當(dāng)單 Channel 數(shù)據(jù)塊小于 NIC 傳輸層采用的 512KiB FIFO 緩沖區(qū)容量時(shí)，將會(huì)發(fā)送未充分填充的緩沖區(qū)，進(jìn)而導(dǎo)致 PCIe 及網(wǎng)絡(luò)吞吐量下降，特別是啟動(dòng)多 QP(Queue Pair)以實(shí)現(xiàn) ECMP(Equal-Cost Multi-Path Routing)負(fù)載均衡的場(chǎng)景。為此，NCCL 采用啟發(fā)式方法動(dòng)態(tài)減小 Channel 數(shù)量(參見 enqueue.cc 中的 calcP2pChunkSize 函數(shù)，這里可能是寫錯(cuò)了，新版本里對(duì)應(yīng) calcP2pChannelCount)。

圖片

NCCL 在 Communicator 初始化階段會(huì)建立一組初始化 Channel，其總數(shù)主要由系統(tǒng)拓?fù)浜图軜?gòu)決定。當(dāng)發(fā)起集合通信操作時(shí)，NCCL 會(huì)動(dòng)態(tài)選擇適合對(duì)應(yīng)任務(wù)的算法和協(xié)議。同時(shí)，NCCL 內(nèi)部調(diào)優(yōu)模型會(huì)根據(jù)所選策略、當(dāng)前消息大小、可用帶寬及每個(gè) Channel 配置的線程數(shù)來確定該操作使用的 Channel 數(shù)量。(PS：NCCL 早期版本支持設(shè)置 NCCL_NTHREADS 等環(huán)境變量來調(diào)優(yōu)通道選擇，但新版本都不推薦此類手動(dòng)調(diào)優(yōu))

每個(gè) Channel 的邏輯通信拓?fù)渲苯佑绊?GPU 間數(shù)據(jù)傳輸方式：

• Ring 拓?fù)渲校焊?GPU 識(shí)別其直接前驅(qū)和后繼節(jié)點(diǎn)以形成單向通信環(huán)。
• Tree 拓?fù)渲校焊?GPU 記錄其父節(jié)點(diǎn)和子節(jié)點(diǎn)的 Rank 號(hào)，構(gòu)建邏輯通信樹。

為提升帶寬利用率，NCCL 采用雙二叉樹結(jié)構(gòu)——兩個(gè)樹中不存在共用的非葉節(jié)點(diǎn)，且至多一個(gè)節(jié)點(diǎn)在兩樹中同為葉節(jié)點(diǎn)。當(dāng)節(jié)點(diǎn)數(shù)為偶數(shù)時(shí)通過鏡像構(gòu)建第二棵樹，奇數(shù)時(shí)則進(jìn)行一個(gè)位置的偏移。這些拓?fù)浣Y(jié)構(gòu)在 Communicator 初始化時(shí)確立，并在所有集合操作中復(fù)用。

對(duì)于使用 ncclGroupStart 和 ncclGroupEnd 的分組 P2P 操作，NCCL 會(huì)盡可能將每次傳輸分配到獨(dú)立 Channel，從而實(shí)現(xiàn)多組獨(dú)立發(fā)送與接收操作的并行執(zhí)行，以此實(shí)現(xiàn)傳輸間的任務(wù)級(jí)并行。

3.4 調(diào)優(yōu)模型

NCCL 調(diào)優(yōu)模型對(duì)應(yīng)的實(shí)現(xiàn)在 tunning.cc 中，對(duì)應(yīng) ncclTopoTuneModel 函數(shù)，其主要功能和流程如下：

線程數(shù)(maxThreads)自動(dòng)推算：

• 根據(jù)算法(如 Ring、Tree、CollNet、NVLS)和協(xié)議(Simple、LL、LL128)，結(jié)合帶寬、通道數(shù)、PCIe帶寬等，自動(dòng)為每種算法/協(xié)議選擇合適的線程數(shù)。
• 支持通過環(huán)境變量(如 NCCL_NTHREADS、NCCL_LL128_NTHREADS)進(jìn)行覆蓋。

硬件特性索引與帶寬上限選擇：

• 根據(jù) GPU 架構(gòu)(Volta/Ampere/Hopper/Blackwell)和節(jié)點(diǎn)數(shù)，選擇合適的帶寬上限表(如下圖所示的 llMaxBws、perChMaxTreeBws 等)。
• 單節(jié)點(diǎn)時(shí)主要看 GPU 類型，多節(jié)點(diǎn)時(shí)還考慮 CPU 廠商(Intel/AMD)。

圖片

帶寬(bandwidth)與延遲(latency)建模：

• 對(duì)每種集合通信操作(如 AllReduce、Broadcast、AllGather、ReduceScatter)、每種算法、每種協(xié)議，計(jì)算理論帶寬和延遲。
• 計(jì)算時(shí)會(huì)考慮：

     a.算法圖參數(shù)(如 bwIntra、bwInter、nChannels)

     b.硬件帶寬上限

     c.節(jié)點(diǎn)數(shù)、每節(jié)點(diǎn) GPU 數(shù)

     d.不同協(xié)議的效率修正(如 LL128 只在特定架構(gòu)/拓?fù)湎聠⒂?

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

     f.網(wǎng)絡(luò)延遲、PCIe/NVLink/網(wǎng)絡(luò)等不同鏈路的延遲模型

算法/協(xié)議使能與用戶控制：

• 支持通過環(huán)境變量 NCCL_ALGO、NCCL_PROTO 精細(xì)控制哪些算法/協(xié)議可用。
• 自動(dòng)禁用不支持的算法(如單節(jié)點(diǎn)禁用 CollNet/NVLS Tree，多節(jié)點(diǎn)無 NVSwitch 禁用 CollNet Direct)。
• LL128 協(xié)議只在特定硬件和配置下默認(rèn)啟用，防止數(shù)據(jù)錯(cuò)誤。

最終調(diào)優(yōu)參數(shù)輸出：

• 計(jì)算并輸出每種算法/協(xié)議/操作的最終帶寬、延遲、線程閾值等參數(shù)，供 NCCL 內(nèi)部調(diào)度和選擇最優(yōu)通信方案。
• 支持通過環(huán)境變量 NCCL_THREAD_THRESHOLDS 覆蓋線程閾值。

四、通信協(xié)議

4.1 概覽

NCCL 采用多種通信協(xié)議以優(yōu)化集合通信的數(shù)據(jù)傳輸效率，如下圖 Table 1 所示，包括 Simple、LL(Low Latency，低延遲)和 LL128，旨在帶寬和延遲之間實(shí)現(xiàn)不同的平衡：

圖片

 4.2 Simple 協(xié)議

Simple 協(xié)議旨在最大化帶寬利用率，適用于大消息傳輸。該協(xié)議通過將數(shù)據(jù)分割為相對(duì)較大的數(shù)據(jù)塊，并通過通信 Channel 進(jìn)行分發(fā)來實(shí)現(xiàn)高效傳輸，能實(shí)現(xiàn)接近峰值的帶寬。

為保持內(nèi)存一致性，采用內(nèi)存屏障來確保數(shù)據(jù)的正確排序與可見性。接收方必須等待完整數(shù)據(jù)塊傳輸完畢后方可訪問數(shù)據(jù)。然而，內(nèi)存屏障會(huì)帶來顯著開銷，對(duì)于小消息傳輸而言，這種同步開銷會(huì)成為主要性能瓶頸，導(dǎo)致其在總體傳輸時(shí)間中占比過大。

4.3 LL 協(xié)議

為解決 Simple 協(xié)議存在的延遲問題，NCCL 引入了針對(duì)小消息傳輸?shù)?LL 協(xié)議。其摒棄了傳統(tǒng)的內(nèi)存屏障機(jī)制，轉(zhuǎn)而采用基于標(biāo)志位的輕量級(jí)同步方案。傳輸數(shù)據(jù)時(shí)伴隨發(fā)送一個(gè)微型標(biāo)志位以標(biāo)識(shí)數(shù)據(jù)有效性，接收端可立即識(shí)別數(shù)據(jù)就緒狀態(tài)，從而避免了高開銷的內(nèi)存屏障操作。

具體來說，LL 協(xié)議采用 8 字節(jié)原子操作實(shí)現(xiàn)數(shù)據(jù)傳輸，每組數(shù)據(jù)包包含 4 字節(jié)有效數(shù)據(jù)和 4 字節(jié)標(biāo)志位。協(xié)議強(qiáng)制要求中間緩沖區(qū)駐留于 Host 內(nèi)存，便于 CPU 通過輪詢標(biāo)志位判斷數(shù)據(jù)是否已就緒并可通過 NIC 發(fā)送。這樣做主要是因?yàn)?PCIe 總線輪詢 GPU 內(nèi)存的延遲遠(yuǎn)高于 DRAM 訪問，且需要顯式同步機(jī)制確保 Host 端數(shù)據(jù)可見性。

該設(shè)計(jì)實(shí)現(xiàn)了低延遲特性，但導(dǎo)致無法使用 GPU RDMA 技術(shù)，嚴(yán)重制約了帶寬性能。實(shí)測(cè)表明，LL 協(xié)議僅能達(dá)到峰值帶寬的 25%-50%(具體數(shù)值取決于互連架構(gòu))。因此，該協(xié)議僅推薦用于時(shí)延敏感而帶寬需求次優(yōu)的小規(guī)模數(shù)據(jù)傳輸場(chǎng)景。

4.4 LL128 協(xié)議

LL128 協(xié)議在保留低延遲特性的基礎(chǔ)上對(duì) LL 協(xié)議進(jìn)行了優(yōu)化，顯著提升帶寬利用效率，尤其是 NVLink 等高性能互連場(chǎng)景。與 LL 協(xié)議類似，同樣采用基于標(biāo)志位的同步機(jī)制以消除內(nèi)存屏障，但將數(shù)據(jù)傳輸單元從 8 字節(jié)擴(kuò)展至 128 字節(jié)。每個(gè)傳輸單元中 120 字節(jié)用于承載數(shù)據(jù)，剩余 8 字節(jié)保留為標(biāo)志位，這使得 LL128 可實(shí)現(xiàn)約 95% 的峰值帶寬利用率。

在網(wǎng)絡(luò)傳輸路徑上，LL128 與 Simple 具有相似性：發(fā)送端 GPU 會(huì)聚合較大數(shù)據(jù)塊后才向 CPU 發(fā)出傳輸就緒通知。雖然這種分塊聚合機(jī)制限制了跨節(jié)點(diǎn)的流水線并行度，但由于其較小的傳輸粒度，LL128 仍能充分利用節(jié)點(diǎn)內(nèi)部的細(xì)粒度流水線優(yōu)勢(shì)。這種低延遲與高吞吐的協(xié)同特性使 LL128 能適配廣泛的消息規(guī)模。

需要注意的是，LL128 協(xié)議對(duì)硬件有更嚴(yán)格要求。其依賴于 128 字節(jié)的原子寫操作，要求內(nèi)存系統(tǒng)或互連架構(gòu)不得對(duì)這類操作進(jìn)行拆分或亂序執(zhí)行。在因 PCIe 限制或其他架構(gòu)約束無法保證原子性的系統(tǒng)中，NCCL 將禁用 LL128 協(xié)議以避免數(shù)據(jù)損壞。因此協(xié)議選擇不僅需考慮消息規(guī)模，還需綜合評(píng)估系統(tǒng)級(jí)能力。

4.4 協(xié)議選擇與對(duì)比

NCCL 在運(yùn)行時(shí)根據(jù)用戶配置(如 NCCL_PROTO 參數(shù))、集合通信算法及內(nèi)部性能啟發(fā)式規(guī)則，動(dòng)態(tài)選擇 Simple、LL 和 LL128 三種協(xié)議。若未顯式指定協(xié)議，系統(tǒng)將基于拓?fù)浣Y(jié)構(gòu)、GPU 架構(gòu)、消息大小等性能指標(biāo)構(gòu)建調(diào)優(yōu)模型，自動(dòng)選擇最優(yōu)的算法-協(xié)議組合。典型場(chǎng)景下，針對(duì)小消息采用 LL/LL128 以降低通信延遲，而對(duì)大消息則選用 Simple 以實(shí)現(xiàn)最大吞吐量。

五、數(shù)據(jù)傳輸策略和傳輸層

5.1 概覽

如下圖 Table II 所示，NCCL 根據(jù)通信發(fā)生在單一節(jié)點(diǎn)(節(jié)點(diǎn)內(nèi)通信)還是跨多個(gè)節(jié)點(diǎn)(節(jié)點(diǎn)間通信)采用了差異化的數(shù)據(jù)傳輸策略與傳輸機(jī)制。每種傳輸機(jī)制均針對(duì)特定硬件架構(gòu)與互連類型進(jìn)行優(yōu)化，以此支撐可擴(kuò)展的集合通信操作。

圖片

5.2 節(jié)點(diǎn)內(nèi)數(shù)據(jù)傳輸

如下圖 Figure 1 所示，NCCL 采用層次化架構(gòu)來實(shí)現(xiàn)節(jié)點(diǎn)內(nèi)通信，優(yōu)先選擇同一物理機(jī)內(nèi) GPU 間延遲最低、帶寬最高的傳輸路徑。該策略深度依賴 NVIDIA 的 GPUDirect P2P 技術(shù)，使得 GPU 能夠直接訪問彼此顯存，無需經(jīng)由 CPU 系統(tǒng)內(nèi)存中轉(zhuǎn)。

圖片

節(jié)點(diǎn)內(nèi)通信核心是 P2P 傳輸層，主要實(shí)現(xiàn)在 src/transport/p2p.cc 文件中。

• 當(dāng) GPU 通過 NVIDIA NVLink 互連時(shí)，NCCL 優(yōu)先采用該路徑，在 NVLink 上實(shí)施 GPUDirect P2P 以利用這些專用的高速直連通道。
• 若 NVLink 不可用，NCCL 仍可通過 PCIe 總線進(jìn)行 GPUDirect P2P 通信，該功能同樣由 P2P 傳輸層管理。該備用方案通常比采用 cudaMemcpy 經(jīng)由 Host 內(nèi)存中轉(zhuǎn)的性能表現(xiàn)更優(yōu)。

NCCL P2P 傳輸中的一項(xiàng)關(guān)鍵優(yōu)化是 P2P_DIRECT 模式，該模式在通信 Rank 屬于同一進(jìn)程時(shí)啟用。雖然單進(jìn)程與多進(jìn)程通信均采用不經(jīng)過 CPU 的 GPU 間直接傳輸，但 P2P_DIRECT 模式通過兩種方式顯著提升效率：

• 首先，在同一地址空間內(nèi)直接使用 GPU 內(nèi)存指針，避免了 IPC 句柄的開銷;
• 更重要的是，該模式采用 directSend/directRecv 等原語消除中間數(shù)據(jù)拷貝，實(shí)現(xiàn)了源緩沖區(qū)和目標(biāo)緩沖區(qū)間的直接傳輸，而非通過 intermediate FIFO 緩沖區(qū)進(jìn)行路由。

在上述優(yōu)化中，NCCL 仍通過共享結(jié)構(gòu)(如 ncclSendMem 與 ncclRecvMem)中的原子 head、tail 計(jì)數(shù)器維持正確同步，確保操作順序性并防止數(shù)據(jù)競(jìng)爭(zhēng)。因此，P2P_DIRECT 在簡(jiǎn)化內(nèi)存尋址與建立更直接傳輸路徑的基礎(chǔ)上，依托 GPUDirect P2P 底層能力實(shí)現(xiàn)了顯著的性能提升。

當(dāng) GPU 間直接 P2P 通信不可用或非最優(yōu)時(shí)，NCCL 會(huì)采用共享內(nèi)存(SHM)傳輸方案。特別是在跨 PCIe 的 P2P 通信中，CPU 對(duì) P2P 數(shù)據(jù)包的處理效率低下會(huì)導(dǎo)致性能下降。SHM 模式通過系統(tǒng)內(nèi)存路由流量，利用 PCIe-內(nèi)存與內(nèi)存-PCIe 傳輸機(jī)制規(guī)避此問題，這種傳輸方式通常能更好地匹配 CPU 的優(yōu)化處理特性。在 SHM 模式下，一個(gè) GPU 的控制進(jìn)程將數(shù)據(jù)寫入共享內(nèi)存段，再由另一個(gè) GPU 的進(jìn)程讀取。

值得注意的是，在某些多 Socket 系統(tǒng)中，若每個(gè) GPU 位于獨(dú)立 CPU Socket 且配備支持GPUDirect RDMA 的本地 NIC 時(shí)，NCCL 會(huì)采用 NIC 實(shí)現(xiàn)節(jié)點(diǎn)內(nèi) GPU 間通信。此時(shí)數(shù)據(jù)將通過 GPU-NIC-NIC-GPU 路徑傳輸，利用 PCIe 帶寬規(guī)避 CPU 互連瓶頸，而非傳統(tǒng) CPU 互連方案。

5.3 節(jié)點(diǎn)間數(shù)據(jù)傳輸

如下圖 Figure 2 所示，NCCL 根據(jù)可用硬件在兩種主要網(wǎng)絡(luò)傳輸協(xié)議之間進(jìn)行選擇：

1.基于套接字的通信：當(dāng) NIC 不支持 RDMA 時(shí)，NCCL 采用套接字傳輸，其實(shí)現(xiàn)位于 transport/net_socket.cc 文件中。在此模式下，中間緩沖區(qū)被分配為 Host 內(nèi)存中的 CUDA 固定內(nèi)存。

• 發(fā)送端首先將數(shù)據(jù)從 GPU 復(fù)制到該緩沖區(qū)，隨后通過標(biāo)準(zhǔn)套接字調(diào)用進(jìn)行網(wǎng)絡(luò)傳輸。
• 接收端則將數(shù)據(jù)接收到主機(jī)緩沖區(qū)后，再將其復(fù)制至 GPU。
• 這種依賴 Host 內(nèi)存作為中轉(zhuǎn)的設(shè)計(jì)，會(huì)導(dǎo)致跨 PCIe 總線的額外內(nèi)存拷貝開銷。
• 發(fā)送和接收過程均遵循 rendezvous 協(xié)議，即在數(shù)據(jù)傳輸發(fā)生前，收發(fā)雙方需協(xié)調(diào)確認(rèn)緩沖區(qū)準(zhǔn)備就緒。

2.IB Verbs 傳輸：針對(duì) IB 或 RoCE 等高性能網(wǎng)絡(luò)，NCCL 采用 net_ib.cc 文件中實(shí)現(xiàn)的 IB 傳輸協(xié)議。該傳輸方案利用 RDMA 技術(shù)，在最小化 CPU 干預(yù)的情況下實(shí)現(xiàn)節(jié)點(diǎn)間的直接數(shù)據(jù)傳輸。與套接字傳輸類似，所有傳輸均通過中間緩沖區(qū)進(jìn)行中轉(zhuǎn)，但該緩沖區(qū)的具體位置取決于硬件支持及配置參數(shù)。

圖片

默認(rèn)情況下，若 NIC 無法直接訪問 GPU 顯存，系統(tǒng)會(huì)在 Host 內(nèi)存中分配中間緩沖區(qū)。

• GPU Kernel 先將數(shù)據(jù)拷貝至該緩沖區(qū)，隨后代理線程發(fā)起 RDMA 寫操作，將數(shù)據(jù)從 Host 內(nèi)存?zhèn)鬏斨吝h(yuǎn)程節(jié)點(diǎn)。
• 接收端流程相反：NIC 將接收到的數(shù)據(jù)寫入 Host 緩沖區(qū)，代理線程協(xié)調(diào)完成從 Host 內(nèi)存到 GPU 顯存的拷貝。
• 代理線程的核心職責(zé)是管理這些 DMA 與 RDMA 操作。與套接字傳輸機(jī)制類似，數(shù)據(jù)傳送前需通過會(huì)合協(xié)議同步發(fā)送方和接收方。

下文重點(diǎn)闡述 IB 傳輸層實(shí)現(xiàn)的關(guān)鍵特性與優(yōu)化策略：

• GPUDirect RDMA 優(yōu)化：該技術(shù)的核心突破在于支持 NIC 直接訪問 GPU 顯存(GDRDMA)，從而消除 Host 內(nèi)存的中轉(zhuǎn)開銷。此優(yōu)化僅在 NIC 與 GPU 掛載于同一 PCIe Switch 時(shí)啟用。此時(shí)中間緩沖區(qū)直接分配于 GPU 顯存，CPU 代理線程通過 nv_peer_mem 或 Linux DMA-BUF 子系統(tǒng)等機(jī)制，向支持 RDMA 的 NIC 注冊(cè) GPU 顯存空間，使 NIC 能直接映射訪問。NIC DMA 引擎可完全繞過 CPU 和 Host 內(nèi)存，直接對(duì) GPU 執(zhí)行 RDMA 讀寫操作。
• Per-peer 多 Channel 連接：為提升帶寬利用率并降低擁塞，IB 傳輸層默認(rèn)每對(duì)遠(yuǎn)程 GPU 與 NIC 建立 2 條邏輯 Channel(由 NCHANNELS_PER_NET_PEER 參數(shù)控制)。每條邏輯 Channel 維護(hù)獨(dú)立的 ncclIbSendComm 結(jié)構(gòu)體，內(nèi)含一組完整的 IB QP 集合。運(yùn)行時(shí)，Host 側(cè)網(wǎng)絡(luò)代理在發(fā)起 ncclNet->isend() 調(diào)用時(shí)交替使用兩個(gè) sendComm 句柄，從而將流量分流至不同 QP 集合。這種輕量級(jí)輪詢策略通過三重機(jī)制提升效能，且不會(huì)引入額外的 GPU 端狀態(tài)維護(hù)開銷：

     a.增大單 QP 的有效數(shù)據(jù)塊大小。

     b.為支持 ECMP 的架構(gòu)提供路徑多樣性。

     c.增強(qiáng)整體互連效率。

• QP 布局設(shè)計(jì)：針對(duì)每對(duì)通信 Rank，RDMA Plugin 創(chuàng)建兩條可靠連接(RC)QP，實(shí)現(xiàn)雙向通道。

     a.forward QP 負(fù)責(zé)大數(shù)據(jù)流傳輸：代理端發(fā)起一個(gè)或多個(gè) RDMA_WRITE 工作請(qǐng)求，將用戶數(shù)據(jù)直接推送至對(duì)端緩沖區(qū)，最終以零字節(jié) RDMA_WRITE_WITH_IMM 操作收尾。該請(qǐng)求的 immediate 數(shù)據(jù)字段編碼了傳輸總量，接收方通過輪詢此字段確認(rèn)傳輸完成。

     b.reverse QP 僅傳輸小的 clear-to-send(CTS)消息，通過單次 RDMA_WRITE 操作通知遠(yuǎn)程緩沖區(qū)地址、rkeys 及標(biāo)簽信息。

     c.雖然理論上可通過單 QP 實(shí)現(xiàn)功能復(fù)用，但將 CTS 隔離至獨(dú)立 Channel 能有效分離延遲敏感的控制流量與帶寬密集型數(shù)據(jù)流，從而最大限度降低網(wǎng)絡(luò)傳輸中的隊(duì)首阻塞(head-of-line blocking)效應(yīng)。

• 基于回環(huán)RDMA_READ的本地刷新機(jī)制：當(dāng)啟用 GPUDirect RDMA 時(shí)，發(fā)送方必須確保所有未完成的 PCIe 寫入操作在 Kernel 處理數(shù)據(jù)前到達(dá) GPU 顯存。NCCL 通過在最后一次接收完成后發(fā)起虛擬 RDMA_READ 實(shí)現(xiàn)此功能。專用"刷新" QP 采用自連接配置，其就緒接收(RTR)階段將本地 QP 編號(hào)作為目標(biāo)地址。該設(shè)計(jì)使得讀取操作始終在 Host 內(nèi)部完成，但 verbs 層仍會(huì)等待先前 PCIe 寫入操作完成，從而以極低開銷實(shí)現(xiàn)存儲(chǔ)排序屏障。

六、NCCL 集合通信算法

6.1 集合通信算法和協(xié)議

集合通信算法是 NCCL 的核心，它實(shí)現(xiàn)了 GPU 間高效、同步的通信。NCCL 通過將每個(gè)集合操作分解為底層通信原語，并將其分配到多個(gè)并行 Channel 中來實(shí)現(xiàn)這些算法。算法選擇取決于具體的集合操作及相關(guān)執(zhí)行參數(shù)，如消息大小和拓?fù)浣Y(jié)構(gòu)。

NCCL 提供 6 種算法，但并非每種算法均適合每種協(xié)議。如下圖 Table III 所示，對(duì)應(yīng) NCCL 2.19 中 5 種集合通信操作支持的算法與通信協(xié)議，該數(shù)據(jù)源自 src/device 目錄下的相應(yīng)頭文件。

• NVLS 與 CollNet 是專為優(yōu)化 AllReduce 性能設(shè)計(jì)的特殊算法，其中 NVLS 還通過利用特定硬件能力支持 ReduceScatter 和 AllGather 操作。
• CollNet 算法適用于網(wǎng)絡(luò)基礎(chǔ)設(shè)施可直接參與集合運(yùn)算的場(chǎng)景，例如采用 NVIDIA SHARP 技術(shù)時(shí)，可將 Reduce 運(yùn)算或其他部分集合計(jì)算卸載至網(wǎng)絡(luò)交換機(jī)執(zhí)行，從而減少數(shù)據(jù)遷移與延遲。
• CollNet 算法依托 NVIDIA SHARP 技術(shù)實(shí)現(xiàn)網(wǎng)絡(luò)輔助的集合操作：

     a.CollNet Direct 支持節(jié)點(diǎn)內(nèi) All2All 通信。

     b.CollNet Chain 則采用線性拓?fù)渑帕?GPU，沿鏈?zhǔn)浇Y(jié)構(gòu)進(jìn)行上行 Reduce 與下行 Broadcast。

• NVLS 算法旨在利用 NVSwitch 的特性，從而提升集合操作效率。標(biāo)準(zhǔn) NVLS 與 NVLS Tree 算法均采用 NVLink SHARP 實(shí)現(xiàn)節(jié)點(diǎn)內(nèi) Reduce，但在跨節(jié)點(diǎn)處理上存在差異：前者通過 CollNet 及支持 SHARP 的交換機(jī)延續(xù) Reduce 過程，后者則采用基于樹狀結(jié)構(gòu)的扇出傳輸。

圖片

6.2 通信原語

NCCL 通過組合一組底層通信原語來實(shí)現(xiàn)高級(jí)集合操作。這些原語構(gòu)成了 NCCL 集合算法的基礎(chǔ)，封裝了跨 GPU 發(fā)送、接收、歸約和復(fù)制數(shù)據(jù)等基本操作。常見原語包括：

• send
• recv
• recvReduceSend
• recvCopySend
• recvReduceCopySend
• 同時(shí)還包括相應(yīng)的 "direct" 變體。

每個(gè)原語代表一種獨(dú)特的數(shù)據(jù)移動(dòng)或計(jì)算模式，其命名規(guī)范清晰體現(xiàn)了操作順序。例如 recvReduceSend 表示 GPU 從對(duì)端接收數(shù)據(jù)，與本地緩沖區(qū)執(zhí)行歸約操作，并將結(jié)果發(fā)送至下一 GPU 的步驟。在執(zhí)行過程中，NCCL 運(yùn)行時(shí)通過循環(huán)步驟迭代調(diào)度這些原語，從而實(shí)現(xiàn)對(duì)不同算法、拓?fù)浣Y(jié)構(gòu)和傳輸層的靈活協(xié)調(diào)。

NCCL 原語的具體行為還受所選通信協(xié)議影響。根據(jù)采用 Simple、LL 或 LL128 協(xié)議的不同，同步機(jī)制、緩沖區(qū)管理和傳輸粒度會(huì)存在差異。需要特別指出的是，這些底層原語針對(duì)源節(jié)點(diǎn)和目標(biāo)節(jié)點(diǎn)數(shù)量固定且較少的集合操作(如 Ring 和 Tree 拓?fù)?，通常涉及單一源?jié)點(diǎn)和目標(biāo)節(jié)點(diǎn)，某些 Tree 拓?fù)渥疃嗳齻€(gè)節(jié)點(diǎn))進(jìn)行了深度優(yōu)化。雖然這種方法使許多標(biāo)準(zhǔn)集合算法能實(shí)現(xiàn)高效運(yùn)行，但對(duì)于 All2All 等需要處理 N 個(gè)源節(jié)點(diǎn)和 N 個(gè)目標(biāo)節(jié)點(diǎn)的通信模式則效率欠佳。

6.3 集合操作的迭代執(zhí)行

NCCL 處理集合操作時(shí)，首先將用戶輸入數(shù)據(jù)分配到可用的通信 Channel 中，從而實(shí)現(xiàn) Channel 級(jí)并行。每個(gè)通道負(fù)責(zé)輸入數(shù)據(jù)的一個(gè)連續(xù)片段，其范圍由元素總數(shù)(count)與 Channel 數(shù)量共同決定。如下圖 Figure 3 所示，數(shù)據(jù)被劃分為若干區(qū)域，各 Channel(如 Channel 0 和 Channel 1)獨(dú)立處理其分配區(qū)段。各 Channel 的起始索引由 workOffset 確定，處理規(guī)模由 channelCount 指定。

圖片

為優(yōu)化數(shù)據(jù)傳輸與計(jì)算效率，NCCL 為每個(gè) Channel 分配固定大小的緩沖區(qū)，其容量取決于所選通信協(xié)議(Simple、LL 或 LL128，詳見下圖 Table IV)。若 Channel 的數(shù)據(jù)區(qū)域超出緩沖區(qū)容量，NCCL 將數(shù)據(jù)分解為若干外層循環(huán)迭代。每次迭代處理不超過緩沖區(qū)大小的數(shù)據(jù)段(每次迭代處理 loopCount 個(gè)元素)，Channel 通過多次循環(huán)完成全部分配元素的處理。

圖片

在每次外層循環(huán)迭代中，NCCL 采用流水線技術(shù)：將 Channel 緩沖區(qū)劃分為固定數(shù)量的 Slot(通常為 8 個(gè)，由 NCCL_STEPS 參數(shù)設(shè)定)。每個(gè) Slot 可獨(dú)立推進(jìn)通信與計(jì)算的不同階段，實(shí)現(xiàn)數(shù)據(jù)傳輸與規(guī)約/復(fù)制操作的流水線重疊。每個(gè)基礎(chǔ)步驟處理一個(gè)數(shù)據(jù)塊(含 chunkCount 個(gè)元素，循環(huán)末塊為 lastChunkCount)，并將其映射至緩沖區(qū) Slot。這種分塊機(jī)制確保通信 Channel 持續(xù)飽和，通過新數(shù)據(jù)塊與進(jìn)行中操作的重疊實(shí)現(xiàn)最大吞吐量。

在 NCCL 中，數(shù)據(jù)移動(dòng)的基本單位稱為元素，其具體含義取決于集合操作類型：

• 對(duì)于 ncclAllGather 和 ncclBroadcast 操作，每個(gè)元素即單個(gè)字節(jié)，因?yàn)檫@些操作的核心目標(biāo)是高效移動(dòng)與拼接數(shù)據(jù)。字節(jié)級(jí)粒度賦予 NCCL 在數(shù)據(jù)打包與傳輸上的靈活性，使其獨(dú)立于底層數(shù)據(jù)類型。
• 對(duì)于 ncclAllReduce、ncclReduceScatter 和 ncclReduce 操作，每個(gè)元素對(duì)應(yīng)著用戶自定義數(shù)據(jù)類型(如 float 或 int)，因?yàn)檫@些運(yùn)算需要基于數(shù)據(jù)類型層面才有意義的算術(shù)規(guī)約操作。

如上圖 Figure 3 展示了該過程的具體實(shí)現(xiàn)。圖中每個(gè)單元格代表 sendBuff 中的一個(gè)數(shù)據(jù)元素。為便于說明，示例設(shè)定 channelCount 為 2，chunkCount 為 2，loopCount 為 4。

• Channel 0 從其 workOffset 起始，以 loopCount 為循環(huán)迭代步長(zhǎng)處理元素，并進(jìn)一步將其分解為 chunkCount 大小的數(shù)據(jù)塊。
• Channel 1 在其對(duì)應(yīng)區(qū)域遵循相同處理邏輯。

6.4 定性算法分析

所有常見的 NCCL 集合通信算法均遵循迭代處理模型，其核心差異在于 GPU 能否對(duì)連續(xù)循環(huán)迭代進(jìn)行流水線化處理。基于這一特性，算法可分為兩類：流水線模式與非流水線模式。

6.4.1 非流水線模式

各 GPU 必須完整執(zhí)行當(dāng)前迭代的所有任務(wù)后，方能啟動(dòng)下一輪迭代。Ring AllReduce、Ring AllGather 及 Ring ReduceScatter 均屬此列。下文分析中，k 表示參與集合通信的 GPU 數(shù)量。

Ring AllReduce：通過分布式規(guī)約階段與數(shù)據(jù)分發(fā)階段的結(jié)合，確保所有 k 個(gè)參與 GPU 均能獲取完整的逐元素規(guī)約結(jié)果。如下圖 Table V 所示，每次循環(huán)包含 2k?1 個(gè)步驟：

圖片

如下圖 Figure 4 所示，Ring AllReduce 可以分為 ReduceScatter 和 AllGather 兩個(gè)階段(PS：我們之前已經(jīng)多次介紹，這里不再贅述)：

圖片

如下圖所示，在 NCCL 中 RingAllReduce 的實(shí)現(xiàn)并不是完全分割成 ReduceScatter 和 AllGather 兩個(gè)操作，而是合在一起通過 2k-1 個(gè) Step 實(shí)現(xiàn)。每個(gè)過程的操作與上圖對(duì)應(yīng)，只是用了 “direct” 方式，比如：

• directSend(Send)
• directRecvReduceDirectSend(RecvReduceSend)
• directRecvReduceCopyDirectSend(RecvReduceCopySend)
• directRecvCopyDirectSend(RecvCopySend)
• directRecv(Recv)

圖片

Ring AllGather：通過 k?1 個(gè)通信步驟，在邏輯 Ring 拓?fù)浣Y(jié)構(gòu)下收集所有 Rank 貢獻(xiàn)的完整數(shù)據(jù)塊集合。每步傳輸中，各 GPU 接收并聚合來自網(wǎng)絡(luò)中其他 Rank 的數(shù)據(jù)塊，最終實(shí)現(xiàn)全局?jǐn)?shù)據(jù)同步。

圖片

其關(guān)鍵階段(如上圖 Table VI 所示)，包括如下圖所示 3 個(gè)主要過程：

• directSend/directCopySend(Send)：若操作為原地執(zhí)行，則該數(shù)據(jù)塊已位于輸出緩沖區(qū);否則，GPU 需通過 copySend 原語將數(shù)據(jù)從輸入緩沖區(qū)復(fù)制至該段。
• directRecvCopyDirectSend(RecvCopySend)：每步接收來自左側(cè)相鄰 Rank 的數(shù)據(jù)塊，將其存儲(chǔ)于輸出緩沖區(qū)的對(duì)應(yīng)段，并轉(zhuǎn)發(fā)至右側(cè)相鄰 Rank。
• directRecv(Recv)：通過 recv 操作接收最后一個(gè)缺失數(shù)據(jù)塊。

圖片

Ring ReduceScatter：對(duì)初始分布于 k 個(gè)GPU上的數(shù)據(jù)塊執(zhí)行逐元素 Reduce 運(yùn)算，隨后將 Reduce 結(jié)果的不同段分散至各 GPU。

圖片

其關(guān)鍵階段(如上圖 Table VI 所示)，包括如下圖所示 3 個(gè)主要過程：

• send：各 GPU 將其一個(gè)本地?cái)?shù)據(jù)塊發(fā)送至相鄰 Rank GPU，啟動(dòng) Ring 數(shù)據(jù)傳輸。
• recvReduceSend：各 GPU 從左側(cè)相鄰 Rank 接收部分歸約數(shù)據(jù)塊，將其與發(fā)送緩沖區(qū)中存儲(chǔ)的對(duì)應(yīng)本地塊逐元素合并，并將新的部分 Reduce 塊轉(zhuǎn)發(fā)至右側(cè)相鄰 Rank GPU。
• recvReduceCopy：各 GPU 接收左側(cè) Rank 的最后一個(gè)數(shù)據(jù)塊，執(zhí)行最終 Reduce 運(yùn)算，并將最終 Reduce 結(jié)果直接寫入其接收緩沖區(qū)。

圖片

6.4.2 流水線模式

NCCL 中的 Tree AllReduce、Ring Broadcast 和 Ring Reduce 算法均采用流水線執(zhí)行模式。

Tree AllReduce：該算法在每次循環(huán)迭代中分為兩個(gè)獨(dú)立階段：Reduce 和 Broadcast。如下圖 Figure 5 所示，以 4 個(gè) GPU 為例。圖中展示的是跨 4 個(gè) Rank 的完整 Tree 結(jié)構(gòu)，但分支結(jié)構(gòu)僅在節(jié)點(diǎn)間建立，節(jié)點(diǎn)內(nèi)以簡(jiǎn)單鏈?zhǔn)竭B接。在 NCCL 的另一種實(shí)現(xiàn)中，這兩個(gè)階段通常將 SM 劃分為兩個(gè)非對(duì)稱組來并發(fā)執(zhí)行：

• 一組負(fù)責(zé)向 Root 節(jié)點(diǎn)執(zhí)行 Reduce。
• 一組同時(shí)執(zhí)行從 Root 節(jié)點(diǎn)開始的 Broadcast。

圖片

Tree AllReduce —— Reduce 階段：

• directSend：葉節(jié)點(diǎn) GPU 通過 send 操作將本地?cái)?shù)據(jù)向上傳遞至父節(jié)點(diǎn)啟動(dòng) Reduce 過程。
• directRecvReduceDirectSend：中間 GPU 使用 recvReduceSend 原語接收子節(jié)點(diǎn)數(shù)據(jù)，與自身數(shù)據(jù)進(jìn)行逐元素 Reduce 后向上傳遞結(jié)果。
• directRecvReduceCopy：最終根節(jié)點(diǎn) GPU 執(zhí)行 recvReduceCopySend 操作，將接收數(shù)據(jù)與本地緩沖區(qū)合并，并將完整 Reduce 結(jié)果復(fù)制至用戶提供的輸出緩沖區(qū)完成 Reduce。

圖片

Tree AllReduce —— Broadcast 階段：

• directSendFromOutput：根節(jié)點(diǎn)通過 recvCopySend 操作將結(jié)果發(fā)送至子節(jié)點(diǎn)。
• directRecvCopyDirectSend：中間 GPU 接收父節(jié)點(diǎn)數(shù)據(jù)后，將其復(fù)制至自身輸出緩沖區(qū)，并使用相同 recvCopySend 原語轉(zhuǎn)發(fā)給子節(jié)點(diǎn)。
• directRecv：葉節(jié)點(diǎn) GPU 則通過簡(jiǎn)單 recv 操作接收數(shù)據(jù)并復(fù)制至輸出緩沖區(qū)。

圖片

其中各類 GPU 角色所使用的設(shè)備原語總結(jié)如下圖 Table VIII 所示。

圖片

Tree Broadcast：將數(shù)據(jù)從用戶指定的根 GPU 傳播至通信組內(nèi)所有其他 GPU。雖然采用 Ring 拓?fù)浣Y(jié)構(gòu)，但其通信模式實(shí)質(zhì)構(gòu)成一條有向鏈，起始于根 GPU 并按順序經(jīng)由各 GPU 傳遞，直至末端 GPU 接收數(shù)據(jù)。

如下圖所示，Tree Broadcast 分為 3 個(gè)階段：

• directSend/directCopySend：

    若根 GPU 的發(fā)送緩沖區(qū)同時(shí)作為接收緩沖區(qū)，則執(zhí)行原地發(fā)送操作;

    否則執(zhí)行 copySend 操作，即先將獨(dú)立發(fā)送緩沖區(qū)的數(shù)據(jù)復(fù)制至接收緩沖區(qū)再進(jìn)行傳輸。

    無論何種情況，根 GPU 都會(huì)將數(shù)據(jù)塊發(fā)送至其 Ring 拓?fù)渲械闹苯雍罄^ GPU。

• directRecvCopyDirectSend：鏈?zhǔn)浇Y(jié)構(gòu)中位于中間的每個(gè)后續(xù) GPU 執(zhí)行 recvCopySend 原語：從其前驅(qū) GPU 接收數(shù)據(jù)塊，將其復(fù)制至自身接收緩沖區(qū)，隨后將數(shù)據(jù)塊轉(zhuǎn)發(fā)至后繼 GPU。此過程持續(xù)進(jìn)行直至數(shù)據(jù)抵達(dá)鏈?zhǔn)浇Y(jié)構(gòu)末端 GPU。
• directRecv：末端 GPU 僅執(zhí)行 recv 操作，將接收數(shù)據(jù)復(fù)制至其接收緩沖區(qū)，無需繼續(xù)發(fā)送，因?yàn)檫壿嬫溨械乃?GPU 此時(shí)均已接收到 Broadcast 數(shù)據(jù)。

圖片

Ring Reduce：對(duì)分布式存儲(chǔ)于多塊 GPU 中的數(shù)據(jù)進(jìn)行逐元素 Reduce 運(yùn)算，最終將聚合結(jié)果匯總至用戶指定的根 GPU 設(shè)備。與 Ring Broadcast 類似，該算法基于 Ring 拓?fù)浣Y(jié)構(gòu)構(gòu)建邏輯傳輸鏈路，數(shù)據(jù)沿該鏈路流動(dòng)并最終向根節(jié)點(diǎn)進(jìn)行累積。

如下圖所示，Ring Reduce 同樣分為 3 個(gè)階段：

• send：算法執(zhí)行時(shí)首塊 GPU(葉子節(jié)點(diǎn))將其本地?cái)?shù)據(jù)塊發(fā)送至環(huán)中的下一節(jié)點(diǎn)。
• recvReduceSend：中間級(jí) GPU 執(zhí)行 recvReduceSend 原子操作：每塊 GPU 接收部分 Reduce 結(jié)果數(shù)據(jù)塊后，與本地對(duì)應(yīng)數(shù)據(jù)執(zhí)行逐元素 Reduce 運(yùn)算，并將更新后的結(jié)果轉(zhuǎn)發(fā)至下一節(jié)點(diǎn)。此過程持續(xù)迭代直至數(shù)據(jù)抵達(dá)目標(biāo)根節(jié)點(diǎn)。
• recvReduceCopy：根節(jié)點(diǎn) GPU 通過 recvReduceCopy 操作完成最終處理：接收最終部分 Reduce 結(jié)果，與本地?cái)?shù)據(jù)進(jìn)行歸約運(yùn)算，并將完全歸約后的輸出存儲(chǔ)至接收緩沖區(qū)。

圖片

6.4 基準(zhǔn)測(cè)試

如下圖 Figure 6 展示了 3 種 NCCL 通信協(xié)議在節(jié)點(diǎn)內(nèi)與節(jié)點(diǎn)間 AllReduce 操作的性能。實(shí)驗(yàn)在瑞士國(guó)家超級(jí)計(jì)算中心(CSCS)的 Alps 超級(jí)計(jì)算系統(tǒng)上進(jìn)行，采用 16 個(gè)配備 NVIDIA GH200 的計(jì)算節(jié)點(diǎn)。每個(gè)節(jié)點(diǎn)提供 150GB/s 的節(jié)點(diǎn)內(nèi)互聯(lián)帶寬，并通過 25GB/s 單向網(wǎng)絡(luò)鏈路接入 Cray Slingshot 互連架構(gòu)。

圖片

在節(jié)點(diǎn)間，對(duì)于 Tree 和 Ring 算法：

• LL 協(xié)議與 LL128 協(xié)議在小消息(小于 64 KiB)場(chǎng)景表現(xiàn)最優(yōu)。當(dāng) AllReduce 消息規(guī)模擴(kuò)展至跨 16 節(jié)點(diǎn)的千兆字節(jié)量級(jí)時(shí)，其性能較 Simple 協(xié)議顯著下降。這主要源于 LL 和 LL128 協(xié)議中基于標(biāo)志位的細(xì)粒度同步開銷，需要通過網(wǎng)絡(luò)處理數(shù)百萬次微小同步操作(每 8 或 128 字節(jié)觸發(fā)一次)。雖然 LL128 憑借更大的緩沖區(qū)尺寸在 NVLink 上展現(xiàn)出極高效率，但 RoCE 網(wǎng)絡(luò)環(huán)境下大規(guī)模節(jié)點(diǎn)間傳輸時(shí)，其累積同步成本超過了這些優(yōu)勢(shì)。在某些情況下，LL128 甚至落后于 LL 協(xié)議，這是因?yàn)殡S著規(guī)模增加，每 128 字節(jié)操作的額外開銷變得顯著，或在激烈競(jìng)爭(zhēng)狀態(tài)下，停頓對(duì)大數(shù)據(jù)單元的影響更為嚴(yán)重。
• 相比之下，Simple 協(xié)議采用更大規(guī)模的傳輸單元和更少的同步事件，使其對(duì)網(wǎng)絡(luò)延遲不敏感，在超大消息傳輸時(shí)能更有效地維持高吞吐量。

另一方面，在節(jié)點(diǎn)內(nèi)通信場(chǎng)景中：

LL128 協(xié)議憑借其充分利用 NVLink 的優(yōu)勢(shì)，在所有消息尺寸下均展現(xiàn)出穩(wěn)定的性能表現(xiàn)。具體而言：

• 對(duì)于小尺寸消息，LL128 性能與 LL 協(xié)議相當(dāng)或僅略遜一籌;
• 在大尺寸消息傳輸時(shí)，其表現(xiàn)幾乎與 Simple 協(xié)議持平。

LL 與 Simple 分別在兩個(gè)極端場(chǎng)景表現(xiàn)最優(yōu)：

• Simple 協(xié)議擅長(zhǎng)處理大消息。
• LL 協(xié)議則專精于小消息傳輸。

最后，可以觀察到，無論在節(jié)點(diǎn)內(nèi)還是節(jié)點(diǎn)間通信環(huán)境下，Ring 算法在大消息傳輸中表現(xiàn)卓越，而 Tree 算法則更適用于小消息場(chǎng)景。

上述基準(zhǔn)測(cè)試可歸納為三項(xiàng)關(guān)鍵發(fā)現(xiàn)：

• 首先，實(shí)驗(yàn)結(jié)果驗(yàn)證了理論預(yù)期——LL 和 LL128 協(xié)議最適合小消息傳輸(尤其在節(jié)點(diǎn)間通信時(shí))，而 Simple 協(xié)議在大規(guī)模分布式傳輸場(chǎng)景中持續(xù)保持性能優(yōu)勢(shì)。
• 其次，必須明確區(qū)分通信環(huán)境屬于節(jié)點(diǎn)內(nèi)還是節(jié)點(diǎn)間，因?yàn)椴煌瑐鬏斔惴?特別是 LL128)在這兩種配置下會(huì)呈現(xiàn)顯著差異的性能特征。
• 最后，雖然人工協(xié)議選擇可用于針對(duì)性調(diào)優(yōu)，但在大多數(shù)情況下依賴 NCCL 的自動(dòng)調(diào)優(yōu)機(jī)制更為有利：允許 NCCL 根據(jù)工作負(fù)載特性自主選擇協(xié)議，通常能在絕大多數(shù)應(yīng)用場(chǎng)景中提供穩(wěn)健的性能與可擴(kuò)展性。

其他測(cè)試如下圖 Figure 7 所示：

圖片

七、參考鏈接

??https://arxiv.org/abs/2507.04786??

??https://github.com/NVIDIA/nccl??

??https://arxiv.org/abs/2505.08936??

本文轉(zhuǎn)載自????AI閑談??????，作者：AI閑談