在過去幾個月中,NVIDIA 集合通信庫(NCCL)開發者一直在努力開發一系列新的庫功能和錯誤修復。在本文中,我們將討論 NCCL 2.22 版本的詳細信息以及解決的痛點。
版本亮點
NVIDIA Magnum IO NCCL 是一個旨在優化 GPU 之間和多節點通信的庫,對于 AI 和 HPC 應用中的高效并行計算至關重要。這次版本的價值在于其新的特點:
- 建立延遲連接以節省 GPU 內存:將連接創建延遲至需要時進行,從而減少 GPU 內存開銷。
- 用于成本估算和工作負載均衡的新 API:提供新 API,幫助您優化計算和通信重疊,或研究 NCCL 成本模型。
- 針對
ncclCommInitRank
的優化和 Instrumentation:消除冗余拓撲查詢,將創建多個 Communicator 的應用的初始化速度提升高達90%。 - 通過 IB 路由器支持多個子網:為跨多個 InfiniBand 子網的作業添加通信支持,從而使 DL 訓練作業能夠在超過 40,000 個端點的 InfiniBand 網絡上運行。
特征
在本節中,我們將深入探討每個新功能的詳細信息:
- 延遲連接建立
- 新的成本模型 API
- 初始化優化和儀器分析
- 新的調諧器插件界面
- 靜態插件鏈接
- 用于中止或銷毀的組語義
- IB 路由器支持
延遲連接建立
NCCL 使用一組靜態分配的持久性連接和緩沖區,以運行其即時數據傳輸協議。對于 NCCL 支持的每種給定算法和協議,它都會創建一組單獨的連接和緩沖區,每個連接和緩沖區都需要數 MB 的 GPU 內存。
作為參考,算法定義了給定集合的參與者之間的高級數據移動,協議定義了 NCCL 發送數據的方式。根據操作、消息大小、規模和拓撲結構,選擇給定的算法和協議,以實現最佳性能。
在 2.22 之前,NCCL 會針對每種組合在對等體之間建立連接,這可能會浪費數 MB 的 GPU 內存來處理永遠無法使用的算法和協議。
現在,NCCL 會等待為給定算法建立連接,直到首次需要它。這將大幅減少 NCCL 內存開銷,尤其是在窄范圍內使用 NCCL 時。例如,如果您反復僅以相同的消息大小運行 ncclAllReduce
,則應在給定系統上僅使用一種算法。
該功能默認處于啟用狀態,但可以通過設置 env NCCL_RUNTIME_CONNECT=0
來禁用它。
在單節點 DGX-H100 的上一個場景中,我們發現僅使用 Ring 算法的 NCCL 的 GPU 內存使用減少了 3.5 倍,僅使用基于 NVSwitch 的歸約時的 GPU 內存使用減少了 1.47 倍。
新的成本模型 API
應用程序開發者希望充分利用NVIDIA系統上提供的計算、內存和帶寬資源。
理想情況下,計算和通信完全重疊,兩者都能完美地完成工作,并將硬件的全部功能發揮到極致。在運行大規模 HPC 和 AI 應用程序時,尤其是在多個平臺上運行一個代碼庫時,很難做到這一點。
為幫助解決此問題,NCCL 添加了一個新的 API,使您能夠了解它認為給定操作需要多長時間。此 API 稱為 ncclGroupSimulateEnd
。其使用方式與 ncclGroupEnd
相同,因此任何熟悉編寫 NCCL 代碼的人都可以輕松使用。
不同之處在于,它不啟動通信操作。相反,NCCL 計算它認為操作需要多長時間,并在提供的 ncclSimInfo_t
結構中設置此操作。
ncclGroupStart() ncclAllReduce() ncclGroupSimulateEnd(sim_t) printf( "Estimated completion time=%f microseconds\n" , sim.time); configureComputeAmount(sim.time, &computeIters, &workItemSize); |
但是,此 API 返回的值并不完全符合現實。這是基于 NCCL 內部模型的估計值。截至 2.22,此 API 僅返回組中最后一次操作的估計時間。
初始化優化和儀器分析
隨著客戶工作負載種類繁多且規模不斷增加,降低 NCCL 初始化的開銷已經成為 NCCL 團隊日益優先考慮的事項。
即使是單節點作業,NVIDIA Hopper GPU 上必須單獨發現和連接的 NVLink 互連數量也有所增加,這導致初始化時間的大幅增加。
我們希望縮短初始化時間,并且必須研究每個初始化步驟的開銷。我們首先分析 ncclCommInitRank
中的每個階段,并研究不同規模的每個階段的時間。現在,每當您收集標準 NCCL 日志(NCCL_DEBUG=INFO
)時,您都會看到這一點。
此外,還有一個新的 NCCL_PROFILE
調試子系統,如果您不關心 NCCL 初始化日志的其余部分,該子系統將僅提供儀器信息。
之前討論過的連接建立是一個前景廣闊的改進領域。切換到 lazy establishment 可以節省內存并縮短初始化時間。
另一個領域是拓撲發現,這是一個初始化步驟,其中每個 NCCL 排名確定節點上可用的硬件。這包括系統上有哪些 GPU 和 NIC、存在多少 NVLink 互連,以及 PCI 拓撲和 NUMA 親和力。
事實證明,NCCL 執行 NVLink 發現的方式并不理想,因為每個等級都在自行發現所有鏈路,從而導致冗余和擁塞。
為解決此問題,我們重復使用了拓撲融合代碼,該代碼最初在 NCCL 2.21 中引入,是多節點 NVLink (MNNVL) 支持的一部分,其中每個節點上的部分可用信息在使用節點間通信的引導過程中進行組合,從而全面了解 NVLink 拓撲結構。
對于 2.22,我們將此功能擴展到在每個節點內運行。現在,每個等級僅會發現有關其自己的 GPU 的信息,然后使用節點內拓撲融合將這些結果與對應的結果相結合。
在單個 8 個 H100 GPU 系統上,延遲連接建立和節點內拓撲融合相結合可以將 ncclCommInitRank
的執行時間縮短 90%(約 6 秒)。以前大約需要 6.7 秒,現在大約需要 0.7 秒。對于在執行期間創建多個通信器的應用程序,這可以大幅減少初始化時間。
新的調諧器插件界面
借助新的調諧器插件接口(v3),NCCL 為插件提供按集合計算的 2D 成本表,報告執行算法和協議的每種組合所需的估計操作時間。
NCCL 將與檢測到的拓撲不兼容的表項設置為 -1
,以便向外部調諧器表明這些組合不支持或不允許被覆蓋。
為選擇特定的組合,外部調諧器會將所需算法或協議組合的值更新為0
或整個表的最小值。在插件更新成本表后,NCCL可以使用它為給定集合選擇最終配置。
靜態插件鏈接
NCCL 團隊提供了一個插件模型,供合作伙伴提供自己的調整或網絡后端,以代替 NCCL 內部模型和 InfiniBand 插件。一些合作伙伴希望以靜態方式將這些插件與應用程序二進制文件關聯起來,以方便起見,并避免加載錯誤的插件。
如果應用程序已靜態鏈接網絡或調優插件,請通過將 NCCL_NET_PLUGIN
或 NCCL_TUNER_PLUGIN
設置為 STATIC_PLUGIN
來指定它。
用于中止或銷毀的組語義
之前,ncclCommDestroy
和 ncclCommAbort
將阻塞調用線程,直到完成。
對于多維并行 ML 工作負載,一個進程管理著多個 NCCL 通信器,并且每個通信器最終必須使用這些 API 拆解。我們為這些應用提供了語義,使其能夠以分組方式一次銷毀多個通信器,從而避免死鎖并提供更好的用戶體驗。
IB 路由器支持
借助此功能,NCCL 可以在由一個或多個路由器連接的不同 InfiniBand 子網中運行。NCCL 會自動檢測兩個通信端點何時位于同一個 InfiniBand 網絡的不同子網上,并交換建立連接和通信所需的 GID 信息。
在子網之間路由時,可使用FLID識別一組要轉發的路由器,并在子網之間實現更高性能和自適應路由。NCCL 2.22將自動檢測是否存在FLID,并將其用于不同子網上端點之間的連接。
問題修復和次要功能
NCCL 2.22 提供以下額外的更新:
- 增加了對 DGX 上的 Google Cloud
allreduce
樹算法的支持。 - 在 IB 異步錯誤中記錄了 NIC 名稱。
- 修復了聚合集合性能問題。
- 修復了注冊發送和接收操作的性能問題。
- 為 NVIDIA 可信計算解決方案添加了基礎架構代碼。
- 為 IB 和 RoCE 控制消息添加了單獨的流量類別,以啟用高級 QoS(使用
NCCL_IB_FIFO_TC
設置)。 - 增加了對分區的Broadcom PCI交換機子部分的PCI對等通信的支持。
總結
NCCL 2.22 版本引入了一些重要的功能和優化,旨在提高高性能計算(HPC)和 AI 應用的性能和效率。改進還包括新的調優插件界面、插件靜態鏈接支持以及增強的組語義,以防止出現死鎖。
有關更多信息,請參閱 Magnum IO 和 NCCL。在 GPU-Accelerated Libraries 論壇上提供反饋。
?