nccl

NVIDIA活跃维护

NVIDIA 官方多 GPU 集合通信库，提供 all-reduce、broadcast 等优化通信原语，是分布式深度学习训练中跨卡数据同步的底层依赖。

Optimized primitives for collective multi-GPU communication

v2.32.6-1发布于 2026-11-09 19:28 UTC5,220star +7 / 7 天1,480fork同步于 2026-11-10 11:36 UTC

多GPU通信集合通信分布式训练 NCCL CUDA

对比 GitHub 官网 communications cpp cuda deep-learning gpu

怎么上手

由 README 提炼，执行前请确认

原文

cd nccl
make -j src.build

硬件要求按 README 说明

README 未说明具体显存要求；需 NVIDIA GPU，支持 PCIe、NVLink、NVSwitch 互联，网络可用 InfiniBand Verbs 或 TCP/IP socket。

需要 CUDA

下载：源码 zip·GitHub Releases（安装包）本站不镜像安装包

项目说明

由 README 自动翻译提炼

没人工点评？写下你的经验

NVIDIA 官方多 GPU 集合通信库，为单机或多机训练提供高带宽通信原语

NVIDIA 官方多 GPU 集合通信库，提供 all-reduce、broadcast 等优化通信原语，是分布式深度学习训练中跨卡数据同步的底层依赖。

核心能力

实现 all-reduce、all-gather、reduce、broadcast、reduce-scatter 等集合通信
支持任意数量 GPU，可跨单节点或多节点部署
针对 PCIe、NVLink、NVSwitch 及 InfiniBand Verbs、TCP/IP 网络优化带宽
支持单进程与多进程（如 MPI）应用调用
可打包为 deb、rpm、txz 及 Python wheel 分发

容易踩的坑

官方推荐直接下载 developer.nvidia.com 的预编译版本，自行编译需先装好 CUDA；CUDA 不在 /usr/local/cuda 时需用 CUDA_HOME 指定路径；测试程序不在本仓库，需另克隆 nvidia/nccl-tests。

内容由模型根据 README 生成，可能不完整或不准确，以官方文档为准。看原文 README

最近发版

5 个版本v2.32.6-1

全部 release

v2.32.6-1NCCL v2.32.6-1 Release最新正式版2026-11-09 · 2026-11-09 19:28 UTC
Vera Rubin Performance
- Adds a pipeline-based cost model with Vera Rubin tuning data for Ring, Tree, NVLS, and NVLSTree.
- Adds Vera Rubin models for symmetric AllGather, ReduceScatter, and AllReduce kernels.
- Speeds up the topology graph search on large NVLink domains such as 32 to 40 GPU VR200 systems, reducing communicator initialization time.
- Tuner plugin interface v7 adds Rubin platform support.
Device API Enhancements
- Adds an experimental Flow device APIs for building pipelined collective kernels over LSA and GIN connections with the Simple, LL, and LL128 protocols. The Flow API signatures are not stable and may change.
- Extends ncclGetPeerPointer and the resource-buffer peer pointer to address every load/store-reachable local peer by world rank, and adds ncclDevCommCanGetPeerPointer to report whether a peer is reachable.
- Adds the NCCL_GIN_CONNECTION_CUSTOM_ARRAY GIN connection type so applications can request connectivity to an explicit list of peers.
- Adds ncclLsaCopyTma, a TMA-based variant of the ncclLsaCopy broadcast on Blackwell and later GPUs.
- Adds a signal-efficient two-signal GIN barrier that GIN plugins can request through the v15 barrierOptions property.
- CFT teams now require and use CUDA logical-endpoint clique information to determine unicast and multicast team sizes.
- Validates GPU and NIC capabilities when establishing GDAKI connections and falls back to the proxy backend when GDAKI is unsupported and NCCL_GIN_TYPE is unset.
- Updates the GPI device backend for the latest libgpi device ABI, including 31-bit signal and counter IDs and extended transfer-size encoding, and updates the bundled GPUNetIO code to v5.0.0-rc3.
- Moves the IR-only C device ABI (ncclGin_C, ncclCoopAny) out of the public device headers into bindings/ir.
Collectives and Runtime Enhancements
- Adds experimental collective kernels called General Kernels (or GENK), built on top of the new Flow APIs. These kernels are compiled out by default to minimize binary size; build with SYM_GENK=1 and set NCCL_SYM_GENK_ENABLE=1 to try them.
- Adds experimental AlltoAll_LL, a latency-oriented symmetric AlltoAll kernel for small and medium per-peer messages on LSA and mixed LSA/GIN communicators, selectable with NCCL_SYM_KERNEL=AlltoAll_LL.
- Supports host RMA on rail-only configurations such as NCCL_CROSS_NIC=0, where host RMA was previously disabled.
- Updates the GIN and RMA plugin interfaces to v15 and v17: connection setup receives an abort flag so ncclCommAbort can interrupt blocking GIN and RMA initialization, and the plugin rankStride is replaced by an explicit peerArray.
- Improves hierarchical CE collective bandwidth when inter-node traffic is distributed across multiple RMA contexts, including under CUDA graph capture.
- Reduces host munmap latency on ATS platforms such as GB200 by granting ncclMemAlloc memory access to the owning GPU only.
Diagnostics and Profiling
- Adds profiler plugin interface v8 with opt-in NCCL-sourced GPU and CPU timestamps on every event.
- Supports loading up to two profiler plugins at a time through a comma-separated NCCL_PROFILER_PLUGIN, each with its own event mask.
- Allows NetPlugin profiler events to be collected independently of proxy events; profiler plugins must accept ProxyStep or NULL parents.
- Updates Inspector to record per-call profiler tags (requiring profiler v7 or later), emit the selected algorithm and protocol per collective in its JSON output, and ship a communication-pattern extractor. NCCL_INSPECTOR_NCCL_SOURCED_TIMINGS defaults to on.
- Adds a RAS diagnostics check verifying that the network devices NCCL selected are consistent across the nodes of a communicator and within each node.
- Adds the DIAG debug subsystem (NCCL_DEBUG_SUBSYS=DIAG) for active and RAS diagnostics details.
- Attaches all GIN/GDAKI queue pairs to a dedicated hardware Q counter set so aggregate count
nccl4py-v0.6.0NCCL4Py v0.6.0 Release2026-09-23 · 2026-09-23 17:51 UTC
查看更新说明
Highlights
- Added experimental CuTe DSL ReduceCopy APIs for LSA, multimem, and local memory operations.
- Added NCCL 2.32 host APIs for collective launch completion events, NVLS configuration, CFT capability inspection, and window registration.
- Added explicit CuTe DSL barrier-session teardown and corrected ThreadScope.THREAD to match libcu++.
New Features

Collective Launch Completion Events
- Added per-collective CUDA launch completion events through NCCLCollConfig.
  
  APIs:
  
  NCCLCollConfig.launch_completion_event
  
  NcclEventSpec
NCCL 2.32 Configuration and Capability Inspection
- Added host-side NVLS configuration, CFT capability reporting, and specialized window registration flags.
  
  APIs:
  
  NCCLConfig.nvls_host_mode
  
  NcclNvlsHostMode
  
  NCCLCommProperties.cft_support
  
  NCCLCommProperties.cft_multicast_support
  
  NCCLCommProperties.cft_counted_support
  
  WindowFlag.GIN_ONLY
  
  WindowFlag.CFT_COUNTED
CuTe DSL ReduceCopy
- Added device APIs for reduction and copy operations across LSA windows, multimem pointers, and local tensors.
  
  APIs:
  
  lsa_reduce_sum(), multimem_reduce_sum()
  
  lsa_copy(), multimem_copy()
  
  lsa_reduce_sum_copy(), multimem_reduce_sum_copy()
  
  local_reduce_sum_copy()
CuTe DSL Barrier Lifecycle
- Added explicit session destruction so barrier handles and indexes can be safely reused.
  
  APIs:
  
  LsaBarrierSession.destroy()
  
  GinBarrierSession.destroy()
  
  BarrierSession.destroy()
Examples and Documentation
- Added 08_reduce_copy.py, demonstrating an LSA reduction across registered windows.
Breaking Changes
- ThreadScope.THREAD now has the libcu++ numeric value 10 instead of 3. Code using the enum member requires no changes; code storing or passing its raw integer value must be updated.
- CuTe DSL barrier sessions must call destroy() exactly once after their final operation. Every thread in the session’s cooperative group must call it from uniform control flow.
Fixes and Enhancements
- The NCCL 2.31.2 device IR issue reported in nccl4py v0.5.0, which could prevent Gin.put() from JIT-compiling with EFA GDA enabled, is resolved in the matching NCCL 2.32.3 IR.
- Added a py.typed marker so compatible static type checkers can use nccl.core annotations.
Compatibility Notes
- The nccl4py 0.6.0 host-side bindings and CuTe DSL API are generated from NCCL 2.32.3 headers.
- When using the CuTe DSL API, use matching NCCL 2.32.3 libnccl.so and libnccl_device.bc; nccl4py does not verify this automatically.
- NCCLCollConfig.launch_completion_event requires a timing-disabled, non-interprocess CUDA event. Every rank must either provide an event or omit it. CUDA 12.3 or later is required for post-launch recording semantics.
- The CuTe DSL device API remains experimental.
v2.32.3-1NCCL v2.32.3-1 Release2026-09-17 · 2026-09-17 21:09 UTC
查看更新说明
Vera Rubin Support
- Adds initial Rubin platform support including support for sm107, CX9 rail and plane detection, and MPS+MLoPart.
- 2.32.3 focuses on new functionalities and does not contain performance model tuning for Rubin. This will be part of the next release to improve out-of-box experience for Rubin users.
Device API Enhancements
- Adds Compute Fabric Transport (CFT) counted-write and wait support.
- Adds socket-based GIN support, enabling custom kernel development with TCP sockets.
- Adds support in GDAKI for LAG-aware QP assignment based on the context ID. (Github PR #2315)
- Adds NCCL_WIN_GIN_ONLY so users can register a window only for GIN usage.
- Optimize GIN performance by skipping mcst operation when possible.
Collectives and Runtime Enhancements
- Adds a ring-based hierarchical copy-engine AllGather implementation selectable with NCCL_HIER_CE_COLL_AG_RAIL_RING_ENABLE. (Github PR #2299)
- Improves Blackwell symmetric AllGather performance, resource overhead modeling and kernel selection with a new cost model.
- Adds optional TLS encryption for NCCL-owned socket traffic when NCCL is built with OpenSSL3, configured through the ncclSetEncryption API.
- Adds ncclCollConfig_t::launchCompletionEvent, allowing callers to observe kernel-launch completion.
- Adds ncclNvlsHostMode_t to ncclConfig_t so applications can disable host NVLS collectives per communicator.
- Optimizes NVLS slot consumption to avoid resource exhaustion issues when using multiple communicators with NVLS.
Diagnostics and Profiling
- Adds the ATTN log level for important non-fatal conditions, such as configuration fallbacks and plugin initialization failures.
- Expands RAS capability with GPU-resident progress counters and a watchdog DMA mirror for diagnosing stalled collective kernels.
- Adds additional RAS diagnostics functionalities for NVLink/NIC state and speed, PCI and GDR configuration, Xid/SXid events and others.
- Reports degraded NVLink fabric bandwidth with an ATTN message during communicator initialization.
- Reports mismatched NCCL Git revisions across communicator ranks during initialization.
Other Improvements
- Fixes PAT+NVLS performance drops on H100 platforms.
- Avoid address space exhaustion when repeatedly registering symmetric windows backed by a single physical memory allocation.
- Reduces communicator initialization overhead at large scale by avoiding scans of inactive proxy poll descriptors.
- Adds an experimental built-in NetworkDirect transport on Windows, with automatic socket fallback.
- EFA team improved EFA GDA support with gin.get API and others. (Github PR #2382)(Github PR #2398)(Github PR #2405)
- Reports Inspector ring-buffer drops and operation counts in JSON. (Github PR #2304)
- NCCL now supports communicators using multiple MIG instances.
- Generates llms.txt for agents to better read NCCL documentation.
Bug Fixes
- Fixes profiler overhead when no profiler plugin is loaded. (Github Issue #2355)
- Fixes P2P IPC registration reuse producing out-of-bounds remote addresses when a registered allocation spans multiple cuMem segments. (Github PR #2362)
- Fixes PAT connection-setup deadlocks when runtime connection is disabled and nodes have uneven local-rank counts. (Github Issue #2385)
- Fixes non-thread-safe token parsing that could corrupt concurrent configuration parsing. (Github Issue #2361)
- Fixes B40 TMA symmetric kernels cra
nccl4py-v0.5.0NCCL4Py v0.5.0 Release2026-09-01 · 2026-09-01 19:41 UTC
查看更新说明
Highlights
- Added host and CuTe DSL APIs for CFT setup and logical-endpoint addressing.
- Added per-collective configuration for algorithm selection, CTA tuning, profiler tags, and vendor options.
- Added communicator capability inspection and support for new NCCL 2.31.2 GIN features.
- Added compile-only CuTe DSL arguments for compiling kernels before creating NCCL resources.
New Features

Per-Collective Configuration
- Added per-call configuration to all collective operations.
  
  APIs:
  
  NCCLCollConfig
  
  VendorOption
  
  The config argument on Communicator.allreduce(), broadcast(), reduce(), allgather(), reduce_scatter(), alltoall(), gather(), and scatter()
Compute Fabric Transport Setup
- Added APIs for configuring CFT resources, querying CFT teams, and resolving logical endpoint addresses. CFT communication and barrier operations are not yet exposed.
  
  Host APIs:
  
  NcclHostCftMode, NcclCftTeamMode, NcclCftCap
  
  NCCLConfig.host_cft_mode
  
  NCCLDevCommRequirements.cft_caps, NCCLDevCommRequirements.cft_barrier_count
  
  Communicator.team_cft(), Communicator.team_cft_multimem
  
  CftLeInfo
  
  RegisteredWindowHandle.get_cft_le_info()
  
  RegisteredWindowHandle.get_peer_le_info()
  
  RegisteredWindowHandle.get_multimem_le_info()
  
  CuTe DSL APIs:
  
  CftTeamMode, CftLeInfo
  
  DevComm.team_cft(), DevComm.team_cft_multimem
  
  Window.cft_le_info(), Window.peer_le_info(), Window.multimem_le_info()
  
  DevComm.resource_buffer_cft_le_info()
  
  DevComm.resource_buffer_peer_le_info()
  
  DevComm.resource_buffer_multimem_le_info()
Communicator and GIN Capabilities
- Added consolidated communicator property inspection and support for new NCCL 2.31.2 configuration and GIN capabilities.
  
  APIs:
  
  NCCLCommProperties, Communicator.properties
  
  NCCLConfig.launch_order_implicit
  
  NCCLConfig.num_rma_sig
  
  NCCLConfig.rma_eager_init
  
  NCCLDevCommRequirements.gin_type
  
  NCCLDevCommRequirements.gin_custom_stride
  
  NcclGinType.EFA_GDA
  
  NcclGinConnectionType.CUSTOM_STRIDE
  
  GinBackendMask.EFA_GDA
Compile Before Creating NCCL Resources
- Added type-only arguments for compiling CuTe DSL kernels before communicators, windows, or device resources are created.
  
  APIs:
  
  nccl.core.device.cute.runtime.make_fake_dev_comm()
  
  nccl.core.device.cute.runtime.make_fake_window()
  
  nccl.core.device.cute.runtime.make_fake_multimem_handle()
  
  nccl.core.device.cute.runtime.make_fake_lsa_barrier_handle()
  
  nccl.core.device.cute.runtime.make_fake_gin_barrier_handle()
Examples and Documentation
- Added a compile-before-resource-creation example and renamed the introductory CuTe DSL example to 00_basic.py.
- See the complete CuTe DSL example suite.
Fixes and Enhancements
- Registered window handles created within nccl.core.group() are now updated correctly after the group completes.
Compatibility Notes
- The nccl4py 0.5.0 host-side bindings and CuTe DSL API are generated from NCCL 2.31.2 headers.
- When using the CuTe DSL API, use the matching NCCL 2.31.2 libnccl.so and libnccl_device.bc; nccl4py does not verify this automatically.
- The CuTe DSL device API remains experimental.
Known Issues
- The NCCL 2.31.2 device IR may fail to JIT-compile Gin.put() when EFA GDA is enabled. Apply this patch to NCCL 2.31.2 and rebuild libnccl_device.bc.
nccl4py-v0.4.1NCCL4Py v0.4.1 Release2026-08-11 · 2026-08-11 23:00 UTC
查看更新说明
Highlights
- Added host APIs for NCCL teams, rank translation, and device resource configuration.
- Expanded the experimental CuTe DSL device API with additional GIN operations, resource addressing, and direct support for host-created resources.
- Added live NCCL parameter access and improved version reporting.
New Features

Teams and Device Resources
- Added APIs for inspecting communicator teams, translating ranks, and configuring resources for device-side communication.
  
  APIs:
  
  NCCLTeam
  
  Communicator.team_world, Communicator.team_lsa, Communicator.team_rail
  
  Communicator.team_rank_to_world(), Communicator.team_rank_to_lsa()
  
  NCCLDevCommRequirements.teams, NCCLDevCommRequirements.resources
  
  NCCLDevCommRequirements.gin_traffic_class
  
  TeamRequirement
  
  LsaBarrierRequirement, GinBarrierRequirement, LLA2ARequirement
  
  MultimemHandle, LsaBarrierHandle, GinBarrierHandle, LLA2AHandle
  
  DevCommResource.multimem_handle(), DevCommResource.resource_handles
  
  RegisteredWindowHandle.get_multimem_device_pointer()
CuTe DSL Device API
- Added device-side rank translation, resource addressing, and additional GIN operations.
  
  APIs:
  
  DevComm.team_rank_to_world(), DevComm.team_rank_to_lsa()
  
  DevComm.resource_buffer_local_pointer()
  
  DevComm.resource_buffer_lsa_pointer()
  
  DevComm.resource_buffer_peer_pointer()
  
  DevComm.resource_buffer_multimem_pointer()
  
  DevComm.resource_buffer_lsa_multimem_pointer()
  
  Window.multimem_pointer(), Window.lsa_multimem_pointer()
  
  Gin.put_value(), Gin.get(), Gin.signal(), Gin.flush()
  
  Gin.read_signal(), Gin.reset_signal(), Gin.signal_shadow_pointer()
  
  Gin.read_counter(), Gin.wait_counter(), Gin.reset_counter()
  
  GIN_ALL_CONTEXTS
  
  ThreadScope, GinResourceSharingMode
  
  DevComm.gin(..., resource_sharing_mode=...)
  
  opt_flags support for Gin.put(), Gin.put_value(), Gin.get(), and Gin.signal()
- Host-created device communicators, registered windows, multimem handles, and LSA/GIN barrier handles can now be passed directly as arguments to @cute.jit functions.
Runtime Inspection
- Added live NCCL parameter access and consolidated version reporting under nccl.core.
  
  APIs:
  
  nccl.core.params
  
  nccl.core.dump_params()
  
  nccl.core.get_version()
  
  nccl.core.show_versions()
Examples and Documentation
- Added CuTe DSL examples covering teams, LSA all-reduce, multimem, GIN operations, barriers, and resource buffers.
Breaking Changes
- nccl.ep is no longer included in nccl4py. The NCCL EP Python bindings are now distributed separately from the NVIDIA/nccl-extensions repository.
- Version helpers are no longer exported from the nccl namespace. Import them from nccl.core.
- Version reporting was consolidated:
  
  nccl.core.get_version() now returns VersionInfo.
  
  nccl.core.Version, get_lib_version(), and get_lib_path() were removed.
  
  VersionInfo.nccl was renamed to VersionInfo.libnccl.
  
  VersionInfo.nccl_ep was removed.
  
  VersionInfo.nccl_bindings was added.
Fixes and Enhancements
- Fixed the default release semantics of Gin.put().
- Improved NCCL parameter error handling and string decoding.
Compatibility Notes
- The top-level nccl Python package is now a PEP 420 namespace package.
- The CuTe DSL device API remains experimental.
- The selected libnccl_device.bc must exactly match the NCCL version used by the bindings; nccl4py does not perform this check automatically.
- When one thread manages multiple local ranks, call Communicator.finalize() within nccl.core.group().
Known Issues
- The following CuTe DSL APIs require NCCL 2.31:
  
  `Gi

用户点评

登录用户的真实使用记录，可按版本和硬件筛选

共 0 条

还没有人点评。第一个说出真实体验的人最有价值。

README 原文

未翻译

在 GitHub 打开

NCCL

Optimized primitives for inter-GPU communication.

Introduction

NCCL (pronounced "Nickel") is a stand-alone library of standard communication routines for GPUs, implementing all-reduce, all-gather, reduce, broadcast, reduce-scatter, as well as any send/receive based communication pattern. It has been optimized to achieve high bandwidth on platforms using PCIe, NVLink, NVswitch, as well as networking using InfiniBand Verbs or TCP/IP sockets. NCCL supports an arbitrary number of GPUs installed in a single node or across multiple nodes, and can be used in either single- or multi-process (e.g., MPI) applications.

For more information on NCCL usage, please refer to the NCCL documentation.

Build

Note: the official and tested builds of NCCL can be downloaded from: https://developer.nvidia.com/nccl. You can skip the following build steps if you choose to use the official builds.

To build the library :

$ cd nccl
$ make -j src.build

If CUDA is not installed in the default /usr/local/cuda path, you can define the CUDA path with :

$ make src.build CUDA_HOME=<path to cuda install>

NCCL will be compiled and installed in build/ unless BUILDDIR is set.

By default, NCCL is compiled for all supported architectures. To accelerate the compilation and reduce the binary size, consider redefining NVCC_GENCODE (defined in makefiles/common.mk) to only include the architecture of the target platform :

$ make -j src.build NVCC_GENCODE="-gencode=arch=compute_90,code=sm_90"

Install

To install NCCL on the system, create a package then install it as root.

Debian/Ubuntu :

$ # Install tools to create debian packages
$ sudo apt install build-essential devscripts debhelper fakeroot
$ # Build NCCL deb package
$ make pkg.debian.build
$ ls build/pkg/deb/

RedHat/CentOS :

$ # Install tools to create rpm packages
$ sudo yum install rpm-build rpmdevtools
$ # Build NCCL rpm package
$ make pkg.redhat.build
$ ls build/pkg/rpm/

OS-agnostic tarball :

$ make pkg.txz.build
$ ls build/pkg/txz/

Python wheel :

$ # Install uv to create the Python wheel (uv manages Python deps in a venv)
$ # See: https://docs.astral.sh/uv/getting-started/installation/
$ curl -LsSf https://astral.sh/uv/install.sh | sh
$ # Build NCCL Python wheel (this also builds the .txz archive as an intermediate)
$ make pkg.python_wheel.build
$ ls build/pkg/python_wheel/

Tests

Tests for NCCL are maintained separately at https://github.com/nvidia/nccl-tests.

$ git clone https://github.com/NVIDIA/nccl-tests.git
$ cd nccl-tests
$ make
$ ./build/all_reduce_perf -b 8 -e 256M -f 2 -g <ngpus>

怎么上手

项目说明

核心能力

容易踩的坑

最近发版

Vera Rubin Performance

Device API Enhancements

Collectives and Runtime Enhancements

Diagnostics and Profiling

Highlights

New Features

Collective Launch Completion Events

NCCL 2.32 Configuration and Capability Inspection

CuTe DSL ReduceCopy

CuTe DSL Barrier Lifecycle

Examples and Documentation

Breaking Changes

Fixes and Enhancements

Compatibility Notes

Vera Rubin Support

Device API Enhancements

Collectives and Runtime Enhancements

Diagnostics and Profiling

Other Improvements

Bug Fixes

Highlights

New Features

Per-Collective Configuration

Compute Fabric Transport Setup

Communicator and GIN Capabilities

Compile Before Creating NCCL Resources

Examples and Documentation

Fixes and Enhancements

Compatibility Notes

Known Issues

Highlights

New Features

Teams and Device Resources

CuTe DSL Device API

Runtime Inspection

Examples and Documentation

Breaking Changes

Fixes and Enhancements

Compatibility Notes

Known Issues

用户点评

README 原文

NCCL

Introduction

Build

Install

Tests

Copyright

相关项目