[SYNC] 2.27.3-1

[rahulvaidya20]

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 Why is it necessary to perform such processing on Intel's CPU when checking bandwidth? NVIDIA/nccl#1162).
• Fix an issue where the network proxy profiler could lose track of an
  event identifier (Issue [nccl-profiling] fix bug on calling wrong id for proxy step in recvProxyProgress NVIDIA/nccl#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 [NCCL2.26.2][RAS] RAS could crash NCCL process during initialization phase NVIDIA/nccl#1669, Potential bug when polling RAS (nccl 2.26.5) NVIDIA/nccl#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 socket: handle short read of the magic NVIDIA/nccl#1613).
• connect() will retry also if poll() did not return a POLLOUT event
  (Issue More logging in socketPollConnect and let some error cases retry NVIDIA/nccl#1618).
• Add error checking in a few instances (Issue Fix logging on connection failure NVIDIA/nccl#1539).
• Fix the loop condition in ncclFindInterfaceMatchSubnet() (Issue Fix the loop condition in ncclFindInterfaceMatchSubnet as it may not behave as expected when maxIfs > 1 NVIDIA/nccl#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 Potential buffer overflow issue in xml reader NVIDIA/nccl#1585).
• Extend logging and state clearing to all four events in the internal IB
  plugin (Issue deprecated NVIDIA/nccl#1650).
• Fix the error path in case IB communication is not ready (Issue Fix comm ready check in ncclIbIsend and ncclIbIrecv NVIDIA/nccl#1489).
• Add ECE logging for IB fabric.
• Fix various minor issues in the graph module (Issue Fix for various issues in the graph module NVIDIA/nccl#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 directSend in treeUpDown need update NVIDIA/nccl#1628).
• Remove duplicate code in sendProxySetup() (Issue delete duplicated code in sendProxySetup() in net.cc NVIDIA/nccl#1420).
• Fix the order of arguments of cudaDeviceCanAccessPeer() (Issue align the peer mem access check with API document NVIDIA/nccl#1507).
• Fix compiler warnings with GCC 14.
• Fix a typo in a comment (Issue Fix typo in channel.cc NVIDIA/nccl#1236).

[rahulvaidya20]

Synced #1794 with develop.

[mustafabar]

Please note that #1861 needs to be merged first. If any concerns, contact me on Teams. CC: @corey-derochie-amd @nileshnegi @wenkaidu

[amd-jnovotny]

@rahulvaidya20 and @corey-derochie-amd Should we add a CHANGELOG entry here?

[corey-derochie-amd]

@rahulvaidya20 and @corey-derochie-amd Should we add a CHANGELOG entry here?

200% yes.