44d92cf9dfea0fab314f0f04f330babdee27a266
11 Коммитов
| Автор | SHA1 | Сообщение | Дата | |
|---|---|---|---|---|
 Kamil Iskra
|
11d4481808 |
NCCL 2.27.3-1
Symmetric memory API and symmetric kernels
* Redesign from the ground up, enabling major latency and bandwidth
improvements.
* Add new API calls to register user-allocated memory among communicator
ranks into a NCCL window: ncclCommWindowRegister() and
ncclCommWindowDeregister(). The calls currently support symmetric
registration for P2P and NVLS, and require VMM memory buffers (i.e.,
CUMEM must be operational).
* Implement specialized kernels taking advantage of symmetrically
registered memory, with performance gains expected particularly for
small to medium message sizes.
* The kernels support 32 bit floating point types and smaller, and sum as
the reduction operator, with no more than one collective operation per
group.
* Floating point summation is always done in fp32 accumulators (with the
exception of fp8 on NVLS, where it uses fp16 inside the switch). Thus,
the accuracy with fp8 and fp16 data types should be much improved.
* This initial implementation supports non-network communicators only (P2P
and NVLS transports).
* To explore this functionality users need to use the new memory
registration API calls with the NCCL_WIN_COLL_SYMMETRIC flag and all
ranks of a communicator must pass buffers at the same offset in the same
registration when invoking a collective NCCL operation.
Add support for DGX Spark.
Add support for DirectNIC (CX8) to the internal IB plugin.
Add a new ncclCommShrink() API call
* It is a non-collective call similar to ncclCommSplit(), which makes it
possible to exclude some (possibly unresponsive) ranks from the parent
communicator.
Add support for loading multiple network plugins
* This enables the creation of generic containers that can work across a
range of providers.
* Allow NCCL_NET_PLUGIN to accept a comma-separated list of plugins to
load.
NVLink SHARP (NVLS) improvements
* Implement NVLS+IB SHARP support for AllGather and ReduceScatter with
user buffer registration. This improves performance and reduces the
number of CTAs needed to achieve peak bandwidth.
* Gracefully fall back by default to other transports if NVLS
initialization fails (the old behavior of returning an error code from a
NCCL call can be preserved by setting NCCL_NVLS_ENABLE=1).
* Decrease the NVLS channel count to 24 on Blackwell systems with multiple
NVLink domains per communicator.
* Enable fine-tuning of NCCL behavior per communicator using new
"ncclConfig_t" members "collnetEnable", "CTAPolicy", and "nvlsCTAs".
Profiler improvements
* Extend the init function by adding communicator name, comm id (hash),
rank, number of ranks, number of nodes, and the NCCL log function to the
argument list. This makes the name and the comm id available to all
events in the communicator without explicitly passing them to each
individual event. Add the communicator id and rank to the profiler trace
filename. Now, the communicator name can be set via a new "ncclConfig_t"
member "commName".
* Improve the accuracy of the GPU kernel events by providing GPU-generated
timestamps for the start and stop of every NCCL operation.
* Harmonize proxy events, removing overlaps between ProxyOp and ProxyStep
states.
* Add support for network-defined event updates (through
"recordEventState").
* Report the correct number of channels used by every collective/p2p
operation (used to be set to nMaxChannels for collectives and absent for
p2ps).
* Fix the logic on proxyCtrl Idle/Active events (Issue #1162).
* Fix an issue where the network proxy profiler could lose track of an
event identifier (Issue #1682).
* Improve the backward compatibility with plugins older than v4.
* Ensure that the work counters are 0-initialized.
* Fix a potential race condition in the network profiler that could result
in an event being linked to a wrong parent.
MNNVL improvements
* Increase to 16 the number of NICs used to communicate between MNNVL
domains on GB200 systems, to optimize the performance of collective
operations.
* Add support for more complex MNNVL topologies with up to 32 NICs per
node.
* If the MNNVL fabric initialization was unsuccessful, NCCL will now fail
by default, so as to avoid inadvertently falling back to a potentially
much slower network transport. Such failures are typically due to a
misconfigured IMEX support on the system. To continue without MNNVL,
restart the job with NCCL_MNNVL_ENABLE=0.
* Fix a potential hang in alltoall-like communication patterns at a scale
of over 80 ranks.
* Make NCCL_P2P_DISABLE=1 imply NCCL_MNNVL_ENABLE=0 (so the latter no
longer needs to be specified on MNNVL systems).
* Fix an initialization failure when NCCL_TOPO_FILE is used on MNNVL
systems.
* Fix the graph search to exclude non-local NICs.
* Fix the SHM transport to use fabric handles on MNNVL systems.
NIC Fusion improvements
* Disable the creation of fused NICs for physical devices that haven't
been merged.
* Flatten multiple ports to a single PCI device within the internal IB
plugin and reparent dual-port NICs under the first PCI parent. If the
parent is not a PCI switch, PCI devices for fused NICs won't be
duplicated.
* Route traffic on GB200-CX8 systems through DirectNIC, not the host
interface.
Improve support for platforms with C2C connectivity (e.g., GB200)
* Enable GPUDirect RDMA for the NICs by default.
* Add support for P2C (PXN over C2C) and the LL128 protocol.
Extend NCCL fault tolerance in multithreaded scenarios
* Support the creation of multiple nonblocking communicators within a
single group and polling in parallel for the completion using multiple
threads (one per communicator).
Enable ncclImplicitOrderLaunch for CUDA 12.9+
* This can potentially speed up NCCL_IMPLICIT_LAUNCH_ORDER.
Improve the netSocket transport latency and control
* Provide finer control over the size of the socket send/receive buffers,
the task size, and the number of sockets that a single peer can open.
* Add support for the inlining of small messages behind the header when
using multiple sockets per connection.
Improve the readability of the CPU affinity in the debug output
* Print it as a range string rather than a bitmask.
Fix a potential race condition in graph execution
* A contention could arise when mixing graph and non-graph execution.
Improve PXN connection code
* Avoid duplicate and unused connections.
RAS fixes
* Fix a memory corruption at job termination time in case of a previously
failed initialization of a RAS socket connection.
* Fix a race condition leading to a crash when generating a RAS report
during communicator initialization (Issues #1669, #1718).
* Fix a potential race condition when gathering data for a RAS status
report.
Fix a potential memory corruption in ncclCommSplit()
* Memory could get corrupted when resource sharing was in use and the size
of the NVLink domain in the new communicator was smaller than in the old
one.
Fix asynchronous graph upload
* Fix a small memory leak.
* Fix oversychronization.
Add a check for out-of-memory conditions in ncclMemAlloc()
Clean up the NCCL socket code
* accept() will retry also if just reading the magic failed (Issue #1613).
* connect() will retry also if poll() did not return a POLLOUT event
(Issue #1618).
* Add error checking in a few instances (Issue #1539).
* Fix the loop condition in ncclFindInterfaceMatchSubnet() (Issue #1574).
* Clean up the debug output, downgrading WARN messages to INFO in
non-critical cases, and printing the peer's address where relevant.
Switch NCCL_DEBUG_FILE to line buffering
* This should help avoid mixed-up partial output lines in multithreaded
cases.
Other minor fixes
* Improve the checks for buffer overflows in the graph code (Issue #1585).
* Extend logging and state clearing to all four events in the internal IB
plugin (Issue #1650).
* Fix the error path in case IB communication is not ready (Issue #1489).
* Add ECE logging for IB fabric.
* Fix various minor issues in the graph module (Issue #1635).
* Clean up the debug output in the graph code, downgrading WARN messages
to INFO in non-critical cases.
* Add a missing argument to a directSend() call (Issue #1628).
* Remove duplicate code in sendProxySetup() (Issue #1420).
* Fix the order of arguments of cudaDeviceCanAccessPeer() (Issue #1507).
* Fix compiler warnings with GCC 14.
* Fix a typo in a comment (Issue #1236).
[ROCm/rccl commit:
|
||
 Sylvain Jeaugey
|
db3bfd118f |
2.24.3-1
Network user buffer support for collectives
* Leverage user buffer registration to achieve zero-copy
inter-node communications for Ring, NVLS and Collnet
Add RAS subsystem
* Create a RAS thread keeping track of all NCCL communicators.
* Add a ncclras tool contacting the RAS thread and getting a
report.
Add fp8 support
* Add support for e5m2 and e4m3 8-bit floating point operations.
* Use Tree/PAT algorithms when possible for better numerical
stability.
Add NIC fusion
* Add a NET API to ask the network plugin to fuse a set of
interfaces together.
* Fuse multiple NICs under the same PCI switch as a single,
larger NIC.
Socket connection failure retry
* Retry in case of socket connection failure (unreachable host)
* Avoid "Software caused connection abort" errors on retries
QP connection failure retry
* Retry in case of IB QP connection failure during ibv_modify_qp.
NET API improvements
* Allow plugins to force a flush in case data and completion
ordering is not guaranteed.
* Indicate when completion is not needed (e.g. for the LL128
protocol), allowing plugins to skip generating a completion.
* Allow for full offload of allgather operations when using one
GPU per node.
NCCL_ALGO/NCCL_PROTO strict enforcement
* Extend NCCL_ALGO/NCCL_PROTO syntax to be able to specify
ALGO/PROTO filters for each collective operation.
* Strictly enforce the ALGO/PROTO filters, no longer fall back
on the ring algorithm when the filtering leaves no option and
error out instead.
Enable CUMEM host allocations
* Use cumem functions for host memory allocation by default.
Improved profiler plugin API
* Avoid dependencies with NCCL includes.
* Add information on whether the buffer is registered or not
Adjust PAT tuning
* Improve transition between PAT and ring at scale.
Fix hangs when running with different CPU architectures
* Detect when we use a mix of GPU architectures
* Ensure Algo/Proto decisions are made based on that unified
state.
Fix FD leak in UDS
* Fix a leak when mapping buffers intra-node with cumem IPCs.
Fix crash when mixing buffer registration and graph buffer registration.
* Separate local and graph registration to avoid crashes when we free
buffers.
Fix user buffer registration with dmabuf
* Make ncclSend/ncclRecv communication with buffer registration functional
on network plugins relying on dmabuf for buffer registration.
Fix crash in IB code caused by uninitialized fields.
Fix non-blocking ncclSend/ncclRecv
* Fix case where ncclSend/ncclRecv would return ncclSuccess in non-blocking
mode even though the operation was not enqueued onto the stream.
* Issue #1495
Various compiler tweaks and fixes
* PR #758
Fix typo in ncclTopoPrintGraph
* Issue #1468
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
5ca1b6c160 |
2.22.3-1
Rework core for NVIDIA Trusted Computing
* Compress work structs so that they are shared between channels
* Utilize the full amount of kernel argument space permitted (4k)
before resorting to work fifo.
* Rework the task preprocessing phase.
* Use a separate abortDevFlag which is kept in sync with abortFlag
using cudaMemcpy operations.
* Rename src/include/align.h to src/include/bitops.h
Add lazy connection establishment for collective operations
* Move buffer allocation and connection establishment to the first
collective operation using that algorithm.
* Accelerate init time and reduce memory usage.
* Avoid allocating NVLS buffers if all calls are registered.
* Compute algo/proto in ncclLaunchCollTasksInfo early on.
* Connect peers in ncclCollPreconnectFunc if not connected already.
* Also move shared buffer creation to the first send/recv call.
Accelerate intra-node NVLink detection
* Make each rank only detect NVLinks attached to its GPU.
* Fuse XMLs to reconstruct the full NVLink topology
Add init profiling to report time spend in different init phases.
* Report timings of bootstrap, allgather, search, connect, etc.
* Add new "PROFILE" category for NCCL_DEBUG_SUBSYS.
Add support for PCI p2p on split PCI switches
* Detect split PCI switches through a kernel module exposing
switch information.
* Update the topology XML and graph to add those inter-switch
connections.
Add cost estimation API
* Add a new ncclGroupEndSimulate primitive to return the estimated
time a group would take.
Net/IB: Add separate traffic class for fifo messages
* Add NCCL_IB_FIFO_TC to control the traffic class of fifo messages
independently from NCCL_IB_TC.
Merges PR #1194
Net/IB: Add support for IB router
* Use flid instead of lid if subnets do not match
* Warn if flid is 0
Optimizations and fixes for device network offload (unpack)
* Double the default number of channels
* Cache netDeviceType
* Fix save/increment head logic to enable Tree support.
Support ncclGroupStart/End for ncclCommAbort/Destroy
* Allow Abort/Destroy to be called within a group when managing
multiple GPUs with a single process.
Improve Tuner API
* Provide to the plugin the original cost table so that the plugin
can leave unknown or disabled algo/proto combinations untouched.
* Remove nvlsSupport and collnetSupport.
Do not print version to stdout when using a debug file
* Also print version from all processes with INFO debug level.
Fixes issue #1271
Fix clang warnings in NVTX headers
* Update NVTX headers to the latest version
Fixes issue #1270
Disable port fusion in heterogeneous systems
* Do not fuse ports if a mix of multi-port and single port are detected.
Fix NVLS graphs search for dual NICs.
* Fix NVLS graph search when we have more than one NIC per GPU.
Fix crash with collnetDirect
* Add separate graph search for collnetDirect, testing alltoall paths
and working similarly to the NVLS search.
Fix hang when nodes have different CPU types
* Add the CPU type to the rank peer info.
* Align all ranks on the CPU type after the first allgather.
* Only use the aligned CPU type for all tuning operations.
Fixes issue #1136
Fixes issue #1184
Fix performance of registered send/recv operations
* Allow for single full size operations
* Add INFO to confirm the registration of send/recv buffers.
Move all sync ops to finalize stage
* Ensure ncclCommDestroy is non-blocking if ncclCommFinalize has
been called.
Improve error reporting during SHM segment creation
Improve support of various compilers
Merges PR #1177
Merges PR #1228
Allow net and tuner plugins to be statically linked
* Search for ncclNet or ncclTuner symbols in the main binary.
Merges PR #979
Plugin examples includes cleanup
* Harmonize err.h and common.h usage.
* Add mixed plugin with both net and tuner.
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
792c33598f |
2.21.5-1
Add support for IB SHARP 1PPN operation with user buffers.
Improve support for MNNVL, add NVLS support and multi-clique support.
* Detect the NVLS clique through NVML
* Exchange XML between peers in the same NVLS clique and fuse XMLs
before creating the topology graph.
* Rework bootstrap allgather algorithms to allow for large allgather
operations intra-node (XML exchange).
Net/IB: add support for dynamic GID detection.
* Automatically select RoCEv2/IPv4 interface by default. Allow to
select IPv6 or even the network/mask.
Reduce NVLS memory usage.
* Add stepSize as property of a connection to allow for different
sizes on different peers; set it to 128K for NVLink SHARP.
Improve tuner loading
* Look for more paths, be more consistent with the network device
plugin.
* Also search for tuner support inside the net plugin.
Improve tuner API
* Add context to support multi-device per process.
Add magic number around comm object to detect comm corruption.
* Add some basic check around communicators so that we can report a
problem when a communicator gets corrupted or a wrong comm pointer
is passed to NCCL.
Fix net/IB error path. Github PR #1164
Fix collnet rail mapping with split comm.
Fix packet reordering issue causing bootstrap mismatch
* Use a different tag in ncclTransportP2pSetup for the connectInfo
exchange and the following barrier.
Fix hang when crossNic is inconsistent between ranks.
Fix minCompCap/maxCompCap computation. Github issue #1184
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
f8886d8687 |
2.12.7-1
Add network communication through another GPU connected with NVLink
(PXN).
Add aggregation of messages coming from different local GPUs through
PXN and going to the same destination.
Add new v5 plugin API with grouped receives and tags.
Add compat for v4 plugins.
Add naming of NCCL threads to help debugging.
Fix NVLink detection and avoid data corruption when some NVLinks are
down.
Add support for Relaxed Ordering for IB.
Add profiling and timing infrastructure.
[ROCm/rccl commit:
|
||
 Ke Wen
|
e51f2a83da |
2.11.4-1
Add new API for creating a reduction operation which multiplies the input by a rank-specific scalar before doing an inter-rank summation (see: ncclRedOpCreatePreMulSum).
Improve CollNet (SHARP) performance of ncclAllReduce when captured in a CUDA Graph via user buffer registration.
Add environment variable NCCL_NET_PLUGIN="<suffix>" to allow user to choose among multiple NCCL net plugins by substituting into "libnccl-net-<suffix>.so".
Fix memory leak of NVB connections.
Fix topology detection of IB Virtual Functions (SR-IOV).
[ROCm/rccl commit:
|
||
|
Ke Wen
|
58a9f19dc3 |
2.10.3-1
Add support for bfloat16.
Add ncclAvg reduction operation.
Improve performance for aggregated operations.
Improve performance for tree.
Improve network error reporting.
Add NCCL_NET parameter to force a specific network.
Add NCCL_IB_QPS_PER_CONNECTION parameter to split IB traffic onto multiple queue pairs.
Fix topology detection error in WSL2.
Fix proxy memory elements affinity (improve alltoall performance).
Fix graph search on cubemesh topologies.
Fix hang in cubemesh during NVB connections.
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
3dd3a1ca66 |
2.9.9-1
Fix crash when setting NCCL_MAX_P2P_NCHANNELS below nchannels.
Fix hang during sendrecv dynamic NVB connection establishment on
cubemesh topologies.
Add environment variable to only use SHARP on communicators beyond
a given number of ranks.
Add debug subsystem to trace memory allocations.
Fix compilation with TRACE=1. (Issue #505)
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
a8908b34ee |
2.8.3-1
Optimization for Tree allreduce on A100.
Improve aggregation performance.
Use shared buffers for inter-node send/recv.
Add NVTX profiling hooks.
Accelerate alltoall connections by merging communication for all
channels.
Add support for one hop communication through NVLink, for faster
send/recv communication on cubemesh topologies like DGX-1.
Improve alltoall scheduling to better balance intra/inter node
communication.
Increase send/recv parallelism by 8x, each warp sending or
receiving to a different peer.
Net: move to v4.
Net: make flush operation asynchronous to accelerate alltoall.
Net: define maximum number of requests.
Fix hang when using LL128 protocol after 2^31 steps.
Fix #379 : topology injection failing when using less GPUs than
described in the XML.
Fix #394 : protocol mismatch causing hangs or crashes when using
one GPU per node.
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
5df2502deb |
Improve robustness of PCI detection
Fallback to default values when class/speed is unknown.
[ROCm/rccl commit:
|
||
|
Sylvain Jeaugey
|
40adc74496 |
2.6.4-1
Add support for network collectives.
Add support for XML topology dump/injection.
Add text values for GDR and P2P Levels, including "NVL".
Add speed detection for PCI, Infiniband and Ethernet cards.
Add CPU detection for ARM and AMD CPUs.
Add support for adaptive routing on Infiniband.
Change NET plugin API to v3 : merge PCI path and GPU pointer
capability into a single structure and add other properties.
[ROCm/rccl commit:
|