* 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 9d72be7b2f.
* 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>
* 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.
* Add support for extended fine-grained system memory pool
* Use hipHostRegisterUncached
* Add "sc0 sc1" flags for LL store on gfx950
* Update after HIP flag is changed to hipExtHostRegisterUncached
This feature tracks the proxy events and status of each send/recv op. ProxyTrace keeps a fixed number of active ops in host mem and dumps the status of each op when the program crashes or hangs.
* First version of new replayer, with comments on future TODOs
* plus minor fixes for UT
* Updated format of recorder, especially in binary department, according to replayer's need
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).
This commit handles DMABUF initialization and call appropriate handling function. This fixes crash in OS with no peermem support and relying on only DMABUF.
* Initial test commit
* Handling Dmabuf_fd opening and closing
* Cleanup
* Use DMABuff or Peermem as needed
* Using user input for ibDmaBufSupportInitOnce
* Revert all changes to rocmwrap.cc
* Revert all changes to rocmwrap.cc
* Changing to func definition braces
* Reverting line removal in utils.h
* useDmaBuf to calculate flushEnabled
Minimize the performance impact of the device kernel profiling support when
the profiler plugin is not loaded.
Reduce the overheads of CUDA graph capturing, which increased in NCCL
2.26.2 for large graphs.
Fix the exchange of enhanced connection establishment (ECE) options to
address potential slowdowns on networks utilizing RoCE.
Test if cuMem host allocations work and if not, disable them. Enabled by
default since NCCL 2.24 if the CUDA driver version is at least 12.6, such
allocations rely on NUMA support, which is by default not available under
Docker. We recommend invoking Docker with "--cap-add SYS_NICE" to enable
it.
Fix an initialization error when running with NCCL_NET_GDR_C2C=1 on
multiple MNNVL domains with non-uniform network configurations across
nodes.
Fix the printing of sub-seconds in the debug log when using a custom
NCCL_DEBUG_TIMESTAMP_FORMAT setting.
w/o the header, experience following errors:
```
error: use of undeclared identifier 'RTLD_NOW'
error: use of undeclared identifier 'RTLD_LOCAL'
error: use of undeclared identifier 'dlerror'
error: use of undeclared identifier 'dlsym'
error: use of undeclared identifier 'dlclose'
```
* removed gfx940 and gfx941
* removed gfx940 and gfx941
* Update "gfx94" to "gfx942" in init.cc
* Updated remaining "gfx94" updates to "gfx942"
* Update filenames and variables from gfx940 to gfx942
---------
Co-authored-by: akolliasAMD <akollias@amd.com>
Profiler improvements
* Add events for CUDA kernel start and end.
* Allow network plugins to generate profiling events
* Enable profiling on a per-operation basis, rather than per-communicator.
* Add support for graph capturing.
Add implicit launch order
* Allow to prevent deadlocks when using multiple NCCL communicators per
device by implicitly ordering NCCL operations using the host program
order. Disabled by default, set NCCL_LAUNCH_ORDER_IMPLICIT=1 to enable.
* Add a complementary mechanism to detect host threads racing to launch
to the same device. Enabled by default, set NCCL_LAUNCH_RACE_FATAL=0 to
disable.
Optimize the PAT algorithm
* Separate the computation and execution of PAT steps on different warps,
allowing to run up to 16 PAT steps in parallel to significantly
accelerate PAT and reduce its linear part.
Add support for setting QoS per communicator
* Add a new trafficClass field to the communicator configuration, to
allow the application to select a particular traffic class for a
given communicator. The meaning of the traffic class is
network-specific and should be set in accordance with the network
configuration.
* For the IB/RoCE plugin, existing config variables such as NCCL_IB_SL
and NCCL_IB_TC take precedence.
Allow to enable GPU Direct RDMA specifically on C2C platforms
* Disabled by default, set NCCL_NET_GDR_C2C=1 to enable.
Do not disable user buffer registration unless PXN is really used
* Only disable UB when a communicator has more than one rank per
node on any node.
RAS subsystem improvements
* Report operation counts separately for each collective operation type.
* Provide details about missing communicator ranks and reliably
distinguish ranks that are no longer a given communicator's members
(now reported as NOCOMM) from those that failed to respond.
Add support for timestamps to NCCL diagnostic messages
* On by default for WARN messages; NCCL_DEBUG_TIMESTAMP_LEVELS can be
used to enable them for other debug levels as well.
* The format can be changed using the NCCL_DEBUG_TIMESTAMP_FORMAT config
variable.
Reduce the memory usage with NVLink SHARP (NVLS)
* Potentially save hundreds of MBs of device memory, considering the
multicast buffer size granularity separately from the address alignment.
Update performance tuning for recent Intel CPUs
* Improve algorithm/protocol selection on recent CPUs such as Emerald
Rapids and Sapphire Rapids.
Improve channel scheduling when mixing LL and Simple operations.
* Make LL operations account for 4x more traffic to ensure LL and simple
operations complete at the same time.
Refactor the plugin code
* Clean up and harmonize the support code across the network, tuner,
and profiler plugins.
Add support for comment lines (starting with #) in the nccl.conf file
* Issue #1540.
Make user buffer registration problems print an INFO instead of a WARN.
Drop support for network plugin interface version 5.
Fix a race condition with split-shared communicators
* NCCL could hang during connection setup if multiple communicators
were grouped together that share resources.
Fix a performance regression when using NCCL_CROSS_NIC=1
* NCCL would unnecessarily alternate rings, breaking the GPU-NIC
associations.
Make GID index detection code more resilient
* Dynamic GID detection code was giving up too soon if the
detected index was not available (e.g., wasn't mapped to the
container's sysfs).
* Issues #1538, #1573.
Fix a race condition with non-blocking operation
* Fix issue when creating a non-blocking communicator after a non-
blocking collective operation on another communicator.
Fix shared memory usage on recent Blackwell GPUs.
* Issues NVIDIA/nccl-tests#287, NVIDIA/nccl-tests#291, #1637.
Fix an error with NIC fusion and IB SHARP when recreating communicators
* Disable the unloading of network plugins
Make the auto-merge failures in the NIC fusion non-fatal
* This could happen when trying to merge IB and RoCE devices.
Fixes to ncclCommAbort
* Fix hangs due to the progress thread spinning indefinitely on the
network progress.
* Reduce the abort time by up to two orders of magnitude.
Fix a crash when libnccl.so was dynamically unloaded
* The RAS subsystem was missing a clean-up handler.
Fix a hang if the network plugin's test() call returns an error.
Fix a hang on heterogeneous architectures
* Ensure we harmonize the tuning to avoid different tuning choices,
causing a hang.
Fix double-free on failed ncclCommInitRank and ncclCommFinalize.
Fix a potential list traversal bug during a group launch of multiple
communicators
* Issue #1599.
Unify the handling of NCCL configuration variables
* Under rare circumstances, some variables specified in the config file
could be ignored.
* misc/msccl: read graphCaptureStatus for every collective call
* fix a bug in checking whether UBR is enabled in MSCCLPP
* cmake: Fix patch reversal order
* misc/msccl: add logging
* Added RCCL env params to control setting the SO_REUSEADDR and SO_LINGER socket options. This can allow control over the number of file descriptors created during bootstrapping.
* Casted the linger value to `int` sooner to avoid a scope of unknown typed-ness.
* Added CHANGELOG entry for this feature.
* ext-src: add MSCCLPP memory registration APIs
* update mem-reg patch with mscclpp helper routine to check if buffer is registered
* RCCL integration of MSCCL++ user-buffer registration APIs
* only include mscclpp_nccl header if ENABLE_MSCCLPP is defined
* ext-src: update mscclpp mem-reg patch
* add helper routine to patch
* check handle before MSCCL++ deregister
* fix typo to replace send buff with recv buff
* in case of no mscclpp registration, dduring deRegister call, ont fall back to rccl deRegister which will return an error
* Apply suggestions from code review
Whitespace suggestions and reducing diffs to avoid future merge conflicts
Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com>
* rename helper functions and change their return type
* set RCCL user-buffer registration to occur if attempting MSCCL++ registration with a buffer in managed memory
---------
Co-authored-by: isaki001 <Ioannis.Sakiotis@amd.com>
Co-authored-by: isaki001 <36317038+isaki001@users.noreply.github.com>
Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com>
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
* Make sure the target device is used for MSCCL
* Enable single process mode by default to use MSCCL in MT
* Create a per-rank state when GPUs share a thread
Add scalable init API
* Add new ncclCommInitRankScalable to allow for passing multiple
unique IDs to the init function.
* Spreads the load onto multiple bootstrap roots, allowing for
constant bootstrap time.
* Requires multiple ranks to create a unique ID, and the CPU-side
ID exchange code to call allgather[v] instead of broadcast.
Accelerate init bootstrap operations
* Reduce the number of calls to allgather.
* Allow roots to reply early to ranks when information is already
available.
* Add an option to use ncclNet instead of sockets to perform
bootstrap allgather operations.
Add PAT algorithms for Allgather and ReduceScatter
* Parallel Aggregated Trees, variation of Bruck algorithm.
* Logarithmic number of network steps for small sizes at scale.
* Only supports one rank per node at the moment.
Add support for registered buffers for intra-node communication.
* Allow registered user buffers to be accessed directly intra-node
* Avoids extra copies in algorithms which permit it, saving
memory bandwidth and helping with compute overlap.
Add profiler plugin API
* New plugin API for profiling
* Supports various levels of profiling, with a hierarchy.
Asynchronous graph allocation
* Make calls to cudaMalloc and cudaMemcpy during graph allocation
asynchronous.
* Significantly speeds up graph capture.
Use fatal IB asynchronous events to stop network operation
* Avoids many other error messages
* Only fatal errors are affected; potentially transient errors
(e.g. port down) do not cause an immediate stop.
Set P2P level to PXB on AMD CPUs when using more than 2 GPUs per node
* P2P would cause a significant performance degradation when using
many GPUs, and therefore many interleaved data flows.
* Disable P2P through the CPU when we have 3+ GPUs per node; keep it
enabled when we only have 2 GPUs.
Improve the init logs to report the real NCCL function.
* Make the log report ncclCommInitRank or ncclCommSplit, rather than
the generic ncclCommInitRankFunc.
Add a parameter to set the location of the user configuration file.
* Add NCCL_CONF_FILE environment variable to set where the user's
configuration file resides.
Increase default IB timeout
* Increase IB timeout value from 18 to 20.
* Should help avoid fatal errors on large RoCE systems.
Add new check for nvidia peermem
* On linux kernels 6.6+, /sys/kernel/mm/memory_peers is no longer
present; check for /sys/module/nvidia_peermem/version instead.
Fix old performance regression when mixing small and large operations.
* Improves distribution of work on channels.
Fix crash when NUMA IDs are equal to -1.
* Can happen when a NIC is a virtual NIC, or when linux doesn't
know which NUMA node a device is attached to
* Issue NVIDIA/nccl-tests#233
Fix tree graph search when NCCL_CROSS_NIC is set to 1.
* Would force NCCL to use the balanced_tree pattern, thereby
disabling LL128 on platforms with 1 GPU+1 NIC per PCI switch.
* Would also try to use alternate rings even though it was not
needed.
Compiler tweaks and fixes
* PR #1177
* PR #1228
Fix stack smash
* PR #1325
Fixes for multi-node NVLink + IB operation
Coverity fixes and comments.
* Moved call to `mscclpp_ncclGetUniqueId` into `ncclCommInitRankFunc` to avoid setting up transport early in environments where MSCCL++ isn't valid.
* Checking `mscclEnabled` for the process and the topology to gate MSCCL++.
* Allowed `mscclForceEnable` to enable MSCCL++.
* Added restrictions around calling MSCCL++ collectives (#1281)
* Added restriction to non-zero 32-byte multiple message sizes to MSCCL++ AllGather.
* Renamed and refactored some mscclpp types.
* Only transmit the MSCCL++ unique id for non-split comm init. For splitting comm, it has already been transmitted. Instead, save the MSCCL++ communicator in child communicators when calling `ncclCommSplit`. Only destroy MSCCL++ communicators when no RCCL communicators remain that use it. Also improved trace logging.
* Disable MSCCL++ when using managed memory buffers as it isn't supported.
* Added datatype and op constraints for MSCCL++ AllReduce.
* Added documentation on MSCCL++ restrictions to the README.
* [BUILD] Support custom CMake flags in MSCCLPP (#1275)
* [BUILD] Support custom CMAKE_PREFIX_PATH in MSCCLPP
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com>
* [BUILD] CMake flags to support build-id in MSCCLPP
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com>
* [BUILD] Fix CMake warnings in MSCCLPP build
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com>
* Wrapped all cmake arguments passed to mscclpp to remove empty arguments and properly format them.
---------
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com>
Co-authored-by: Corey Derochie <corey.derochie@amd.com>
* Link to libmscclpp_nccl statically (#1282)
* Switched mscclpp_nccl to static linking. Added a build step to rename the NCCL API functions.
* Undid separation of building libmscclpp_nccl from building librccl with MSCCL++ integration. With a static build, it's either fully enabled or fully disabled.
* `nm` isn't always available in docker containers due to being stripped down. Removed use of `nm` in `cmake` and hard-coded the output into mscclpp_nccl_syms.txt.
* Removed IBVerbs dependency for integrating with MSCCL++ (#1313)
* Renamed `RCCL_ENABLE_MSCCLPP` to `RCCL_MSCCLPP_ENABLE` to conform to MSCCL. Set `RCCL_MSCCLPP_ENABLE` to 1 by default if `ENABLE_MSCCLPP` is defined, or 0 otherwise. Added a log warning if `RCCL_MSCCLPP_ENABLE` is set to 1 but `ENABLE_MSCCLPP` is not defined. (#1294)
* Include mscclpp as a git submodule (#1314)
* Added the desired mscclpp commit as a git submodule.
* Added step to automatically checkout the mscclpp submodule if it isn't already present, in case the user forgot to clone recursively.
* Added instruction to README to clone using --recurse-submodules to get the mscclpp submodule.
* Enabled MSCCL++ feature build.
---------
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com>
Co-authored-by: Nilesh M Negi <Nilesh.Negi@amd.com>
* Added `RCCL_MSCCL_ENABLE_SINGLE_PROCESS` runtime flag to return to the original MSCCL enablement behaviour except when explicitly enabling for multi-thread.
* Added documentation for the new `RCCL_MSCCL_ENABLE_SINGLE_PROCESS` runtime env var.
* adding all nccl apis to api_support to enable rccl tracing by rocprofv3
Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com>
Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
BertanDogancay
Avinash
ycui1984
Arm Patinyasakdikul
Wenkai Du
Rahul Vaidya
Atul Kulkarni
Nilesh M Negi
Dingming Wu
Tim
corey-derochie-amd
Kamil Iskra
Dingming Wu
Nusrat Islam
Pedram Alizadeh
Bertan Dogancay
isaki001
Sylvain Jeaugey
Ziyue Yang
Mustafa Abduljabbar
Mustafa Abduljabbar
mberenjk