Device API (Experimental)
* Introduces device-side APIs to integrate NCCL communication directly into application kernels.
* Supports LSA (Load/Store Access) for CUDA P2P communication over NVLink and some PCIe platforms.
* Supports Multimem for hardware multicast using NVLink SHARP.
* Adds initial framework for GIN (GPU-Initiated Networking), currently under development.
* Introduces device communicators created using ncclDevCommCreate.
* Enables device-side communication operations with synchronization (ncclLsaBarrierSession) and memory accessors (ncclGetLsaPointer, ncclGetLsaMultimemPointer).
* Experimental APIs - signatures and functionality may evolve in future releases.
* No ABI compatibility is guaranteed — applications must be recompiled with each new NCCL release.
Symmetric memory improvements
* Support for aggregating symmetric operations using ncclGroupStart/End APIs.
* Reimplement symmetric kernels using device API.
New Host APIs
* Introduce new host collective APIs: ncclAlltoAll, ncclScatter, ncclGather.
CE (Copy Engine) Collectives
* Reduce SM utilization for alltoall, scatter, gather, and allgather within a single (MN)NVL domain.
* Free up SM capacity for the application to do computation at the same time.
* To enable the feature for ncclAllGather, ncclAlltoAll, ncclGather, ncclScatter, register buffers into symmetric windows and use the NCCL_CTA_POLICY_ZERO flag in the communicator config_t.
NCCL Inspector Plugin
* Introduces an Inspector plugin for always-on performance monitoring.
* Produces structured JSON output with metadata, execution time, bandwidth, and optional event traces for each NCCL operation.
* Enables integration with analysis tools such as Performance Exporter to visualize NCCL performance bottlenecks.
* Lightweight to enable via environment variables NCCL_PROFILER_PLUGIN and NCCL_INSPECTOR_ENABLE.
CMake support (Experiemental)
* Adds a CMake build system as an alternative to existing Makefiles.
* Known issues: pkg.build and Device API currently do not work with CMake.
* The known issues will be addressed in a future release.
Decreased max CTA count from 32 to 16 on Blackwell
* SM overhead is decreased by 50% with this improvement.
* This may cause some perf drop on Blackwell because of the reduced SM usage.
* If the extra SM capacity is not desired, two options are available to restore to previous behavior: 1) Setting NCCL_MIN_CTAS=32 NCCL_MAX_CTAS=32 environment variables; 2) setting communicator config to over-write max CTA count to 32.
* Based on community feedback, future versions may consider different trade-offs between performance and SM overhead.
Plugins
* Network
* App-aware Network plugin. NCCL passes information about communication operations to be executed on the network end point. This allows for better tuning of network end points and their use in the plugins.
* Improve handling of physical and virtual network devices and load/unload.
* Network plugin version 11 - add explicit context and communication ID support for per communicator init/finalize.
* Add Multi-Request Net API. Using this will help NCCL to anticipate multiple send/recv requests and optimize for it. See maxMultiRequestSize field in ncclNetProperties_v11_t.
* Profiler
* Add support for API events (group, collective, and p2p) and for tracking kernel launches in the profiler plugin.
* Add Inspector Profiler Plugin (see section above).
* Add a hook to Google’s CoMMA profiler on github.
* Tuner
* Expose NCCL tuning constants at tuner initialization via ncclTunerConstants_v5_t.
* Add NVL Domain Information API.
* Support multiple plugin types from a single shared object.
New Parameterization and ncclConfig changes:
* Add new option NCCL_MNNVL_CLIQUE_ID=-2 which will use rack serial number to partition the MNNVL clique. This will limit NVLink domains to GPUs within a single rack.
* Add NCCL_NETDEVS_POLICY to control how NET devices are assigned to GPUs. The default (AUTO) is the policy used in previous versions.
* Add NCCL_SINGLE_PROC_MEM_REG_ENABLE control variable to enable NVLS UB registration in the “one process, multiple ranks” case as opt in.
* Move nChannelsPerNetPeer into ncclConfig. NCCL_NCHANNELS_PER_NET_PEER can override the value in ncclConfig.
* Enable PxN over C2C by default
* PxN over C2C will improve performance for Grace-Blackwell platforms by allowing NCCL to leverage the NIC attached to a peer GPU over NVLINK, C2C, and PCIe.
* This behavior can be overridden by setting NCCL_PXN_C2C=0.
Other Improvements:
* Allow FP8 support for non-reductive operations on pre sm90 devices. (See https://github.com/pytorch/pytorch/pull/151594#discussion_r2135777776)
* Fix NVLS+CollNet and temporarily disables COLLNET_CHAIN for >8 GPUs.
* Only consider running interfaces for socket traffic. NCCL will not attempt to use interfaces that do not have the IFF_RUNNING bit. (https://github.com/NVIDIA/nccl/issues/1798)
* Modernize mutex management. Convert to std::mutex and std::lock_guard.
* Remove sm35 and sm50 GENCODE targets which have long been deprecated and were causing issues with the latest NCCL release builds.
* Improved NVLS/NVLSTree tuning prediction to improve algorithm and protocol selection.
* NVLSTree Tuning Fixes. Update tuning data for H100, GB200-NV72.
* Respond better to RoCE link flaps. Instead of reporting an “unknown event” it will now report “GID table changed”.
* Move libvirt bridge interface to the end of possible interfaces so that they are considered last. These interfaces are usually virtual bridges to relay traffic to containers running on the host and cannot be used for traffic to a remote node and are therefore unsuitable.
[ROCm/rccl commit: f1308997d0]
* [rocm_regression] Return errors when HSA_NO_SCRATCH_RECLAIM=1 even for rocm >= 6.4.0
* [rocm_regression] Check firmware version
* [rocm_regression] Resolve review comments
* [rocm_regression] Move hsa env checking into init once func
* [rocm_regression] Prevent hot fix version in firmware
* [rocm_regression] Improve unit tests
[ROCm/rccl commit: 361d596229]
* Support pipelining codegen and template specialization
* Support ReduceCopy pipelining for AllReduce, ReduceScatter, and Reduce (currently enabled for bfloat16)
* Remove need for FUNC_INDEX_TOTAL
* Add pipeline field to device function key construction logic
* Avoid unneeded codegen for LL/LL64 kernels
* Modify conditions and add pipeline dtypes env
* Optimize selection for both gfx942 and gfx950
* Increase pipeline bitfield width
* Use __forceinline__ for all device functions
* Realign reduceCopy with original form
* Add opt-out option to enable perf debugs
* Remove force-reduce-pipelining option from README
* Update CHANGELOG.md
---------
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
[ROCm/rccl commit: 277747c199]
* add direct allgather algorithm
* minor fix
* add debug print for memory allocation tracker
* add message size threshold for direct allgather
* scatter transfers across ranks
* update changelog
* minor fix
* Update CHANGELOG.md
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
* enable direct AG when pxn is ON on MI300X or MI350
---------
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
[ROCm/rccl commit: 5e7937effb]
- Updated ncclDevFuncId to use a hash-based lookup with std::unordered_map.
- Keys are now 64-bit integers, which pack coll, algo, proto, devRedOp, and type fields.
- Improved flexibility and maintainability by moving away from row-based indexing.
- Added error handling for missing keys in the hash map.
- Aligned key generation logic with generate.py and updated generate.py.
[ROCm/rccl commit: c1b3cd8911]
* change gfx950 algo/proto selection for multinode allreduce, allgather, reduceScatter
* gfx950 tuning: enable tuning for broadcast, allreduce starts LL128 earlier and switches to ring earlier, change LL128 start for allgather and reduceScatter
* lower LL128 threshold
* update reduceScatter LL128 min to match LL max for consistency
* enable multinode PXN and increase chunksize for gfx950
* change LL128 start to 128KB, adjust ring-start according to node-count
* disable code-path for fused-AR on LL128 for gfx950
* use LL128 starting from 1KB for multinode allgather on gfx950
* start LL128 earlier for multinode reduceScatter on gfx950
* start LL128 earlier for multinode broadcast on gfx950
* set multinode allreduce to start simple on 64MB for gfx950
* start LL128 from 1KB for multinode broadcast on gfx950
* setting multinode AR to use tree instead of ring at 16MB, 64MB, 128MB
* set multinode broadcast to use LL for up to 256KB depending on node-count for gfx950
* adjust algo for 32MB multinode allreduce on gfx950
* make 32MB tree LL128 for multinode AR on gfx950
* make sure ring is not picked on 2N allreduce on small sizes
[ROCm/rccl commit: 44121db890]
Leverages the traits of extended-scope fine-grain memory to get rid of a device-scope acquire-release fence. This improves throughput for single node workloads on gfx942 and gfx950 for some input sizes (e.g., ~32 MiB to about 256 MiB) when using the simple protocol. Multinode workloads on MI300X see a smaller but statistically significant uplift for some message sizes. Runtime disablement is supported via setting the environment variable RCCL_GFX942_CHEAP_FENCE_ON to 0.
[ROCm/rccl commit: 1aa2570b48]
* Added useAcc as a template parameter to address the 2% performance regression in allreduceWithBias
---------
Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com>
[ROCm/rccl commit: c61152baa4]
- Introduced double-buffering to reduce copy overhead and overlap BF16 arithmetic with data prefetching.
- Aimed to improve performance of reduction-based collectives by up to 10%.
- Implemented based on recommendations from Guennadi Riguer (AMD)
- Added --force-reduce-pipeline option to install.sh to activate this optimization for BF16 reductions.
- Feature is disabled by default to prevent regressions with large messages until auto-tuning logic is upstreamed.
---------
Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
Co-authored-by: Pedram Alizadeh <pmohamma@amd.com>
[ROCm/rccl commit: 0ce20e7e07]
Prevent initialization failures in certain configurations when attempting
to load fp8-specific symmetric multicast kernels on GPUs older than
Blackwell.
[ROCm/rccl commit: 593de54e52]
* Support fused all reduce and elementwise operations
Add additional "acc" parameter to RCCL Replayer logs
Add flag which indicates availability of new API
* Fix Recorder json parsing
* Remove unreachable code
* Remove extra acc pointer check
* .
* Revert "[DEVICE] Adding ability to choose unroll factor at runtime (#1734)"
This reverts commit 4cadf3597c.
* Use noinline to reduce kernels linking time
* Don't use noinline for gfx942 and gfx950 to avoid perf regression
---------
Co-authored-by: AtlantaPepsi <timhu102@amd.com>
Co-authored-by: BertanDogancay <bertan.dogancay@gmail.com>
[ROCm/rccl commit: 9a4213356d]
Boosts single node bfloat16 allreduce performance by up to 20% for some data sizes and provides gating with the RCCL_GFX942_CHEAP_FENCE_OFF environment variable
[ROCm/rccl commit: 11fabf1de1]
* Increased max stack size to 640
* Added new binary for executing unit tests
Added new unit tests for argcheck.cc and alt_rsmi.cc files
Modified the method to execute unit tests to cover static methods
by using a bash script to convert static to non-static functions
and variables on the fly restricted to debug build type.
[ROCm/rccl commit: 275fdd43c1]
Improve support for DirectNIC (CX8)
* Add support for XDR speed detection.
* When DirectNIC is enabled, report only the RDMA interfaces.
Extend the P2C (PXN over C2C) support to send/receive operations.
Support compilation with GCC 14 (Issues #1743, #1751).
Fix the unloading of network plugins that also provide tuner capability.
Fix the change of the current device across the calls to ncclCommDestroy()
and ncclCommAbort().
A note for users on MNNVL systems: please ensure an adequate stack size for
NCCL threads. While the default Linux stack size limit of 8192 KB is known
to be sufficient, we've seen crashes if the limit is changed to
"unlimited", as it causes the glibc library to unexpectedly *decrease* the
stack size of NCCL's background threads to just 2048 KB. Use "ulimit -s"
in bash to print the current limit; if needed, reset it to 8192 KB using
"ulimit -s 8192" (one also needs to ensure that the new setting is
propagated to other nodes when launching a multi-node NCCL job).
[ROCm/rccl commit: 7c12c627c6]
* Adjust tuning factor impacting more than 2 nodes
* Scale max LL128 size for > 2 nodes
* Retune max LL128 range for N > 2
[ROCm/rccl commit: 058264b3f3]
* removing extra build time by removing the gfx11xx arch from using hip_fp8
---------
Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com>
[ROCm/rccl commit: 697bee4ee8]
Wenkai Du
isaki001
Mustafa Abduljabbar
Mark Santesson
ycui1984
BertanDogancay
Nusrat Islam
Nilesh M Negi
Arm Patinyasakdikul
alex-breslow-amd
mberenjk
Karthikeyan Arumugam
Avinash
Sarat Kamisetty
Atul Kulkarni
Kamil Iskra
Rahul Vaidya
Bertan Dogancay