Optimized primitives for inter-GPU communication.
This branch carries AI4SAI changes for SAI UltraPOD and SlimPOD GPU fabrics.
The first optimized API is ncclAlltoAll(): eligible small messages use a
two-stage GPU-island aggregation path, while eligible large messages use a
phased P2P planner. NCCL-SAI also has a narrowly guarded local P2P
transport-selection path that can affect any operation on an eligible
single-host communicator. The validated dual-rail mode requires
NCCL_SAI_FABRIC_PROFILE=ultrapod-fullmesh, NCCL_IB_MERGE_NICS=0, and
NCCL_CROSS_NIC=0. GPUs with exactly two topology-local network devices then
use channel parity, and ordinary ring/tree graph endpoints use the corresponding
local device; CollNet and NVLS endpoint selection remains unchanged. Other
profiles and ineligible layouts retain the corresponding upstream path.
Collective algorithms remain unchanged unless explicitly documented and
validated. See
docs/sai/README.md and docs/sai/COMMUNICATION_TUNING_MATRIX.md for scope,
rollback knobs, and validation requirements.
NCCL-SAI modifications are maintained by AI4SAI/SAI contributors. This project
is derived from NVIDIA NCCL and is not endorsed by NVIDIA. Original NVIDIA NCCL
copyright and license notices are retained; see LICENSE.txt and
docs/sai/NOTICE.md.
For SAI users, the intended runtime mode is drop-in replacement: put the
NCCL-SAI build's lib/ directory before the system NCCL in LD_LIBRARY_PATH.
SAI site modules or prologs can enable transparent SAI behavior by setting
NCCL_SAI_FABRIC_PROFILE to a recognized product-family profile. The current
island and phased AlltoAll defaults are selected by ultrapod-fullmesh. The
dual-rail behavior described above additionally requires explicit merge- and
cross-NIC settings. Broader family names such as ultrapod and slimpod are
recognized activation namespaces but do not imply that the same
topology-specific behavior is valid for every layout.
Unknown profile names and unsupported layouts fall back to upstream NCCL
behavior unless an expert explicitly opts in with NCCL_SAI_A2A_ENABLE=1 and
the related controls.
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.
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.buildIf 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_70,code=sm_70"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/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>Original NVIDIA NCCL source code and documentation retain their upstream
copyright notices. NCCL-SAI modifications are copyright (c) 2026, AI4SAI
contributors and are redistributed under LICENSE.txt; see
docs/sai/NOTICE.md for the modification boundary.