Commit Graph

79 Commits

Author SHA1 Message Date
Wenkai Du 4c9c1d41ee Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 9a077e6947]
2022-11-03 21:17:42 +00:00
Wenkai Du 4630365f2d Fix P2P scheduling
[ROCm/rccl commit: 72ef100050]
2022-10-31 08:54:34 -07:00
Sylvain Jeaugey 0b20e8b7e9 2.15.5-1
Fix crash with CollnetChain on some node topologies
Fix hang when interleaving the capture of different graphs
Fix hang during init in multi-threaded mode
Fix potential data corruption with LL128 protocol on unaligned buffers.
Fix CPU usage during preconnect
Fixes double-free in the error path for ncclCommInitAll
Workaround hang on H100 with Ring/LL128 on 2 GPUs.


[ROCm/rccl commit: cb111f764a]
2022-10-25 00:55:55 -07:00
Wenkai Du 36e5e02e46 Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 4f0e223db4]
2022-10-20 15:41:29 +00:00
Edgar Gabriel f2736a4fb3 introduce a hw topology aware bintree
for hayabusa architecture.


[ROCm/rccl commit: e645b02cd8]
2022-10-03 15:26:21 +00:00
Wenkai Du f6da79844a Add LL128 tuning (#630)
[ROCm/rccl commit: 021932b3c8]
2022-09-27 09:39:09 -07:00
Sylvain Jeaugey b4bac0d15a 2.15.1-1
Add support for H100 (sm90).
Make sure NCCL kernel honor user stream priorities.


[ROCm/rccl commit: da8152e57a]
2022-09-27 02:31:13 -07:00
Edgar Gabriel 95d6ed2154 make binary tree work on 2.13.4
[ROCm/rccl commit: 8f3219dbd4]
2022-09-15 00:01:54 +00:00
Edgar Gabriel 4c17f4dcc1 Merge branch 'develop' into 2.13.4
[ROCm/rccl commit: be935d7ce7]
2022-09-13 17:19:04 -05:00
Edgar Gabriel 7148c0aa7b add binary tree
In addition, introduce the ability to have 2 trees at the same time.
Only for allreduce at the moment.


[ROCm/rccl commit: 65e2ae20e5]
2022-09-13 20:52:32 +00:00
Wenkai Du 7874a99c75 Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: a79d9e3586]
2022-09-09 16:05:38 +00:00
Wenkai Du fe99249cde Enable LL128 protocol support (#605)
* Enable LL128 protocol support

* Use shared memory object directly when possible

[ROCm/rccl commit: 7bbce085cc]
2022-09-08 14:45:27 -07:00
gilbertlee-amd 616cb39a0b Adding opt-in hipGraph support for RCCL via RCCL_ENABLE_HIPGRAPH (#608)
Adding opt-in hipGraph support via RCCL_ENABLE_HIPGRAPH

[ROCm/rccl commit: 47b2fc3a30]
2022-09-06 10:29:46 -06:00
Wenkai Du f18868f439 Use hipExtLaunchKernel when not using graph and not in group mode (#606)
[ROCm/rccl commit: c9f2fe1f65]
2022-08-26 13:40:37 -07:00
akolliasAMD 1d55fe756c Simple tree changes (#599)
changed treebase to create basic balanced tree

[ROCm/rccl commit: 3c1b1ec8c8]
2022-08-19 13:51:49 -06:00
Sylvain Jeaugey f6e1e3d9ed 2.14.3-1
Add support for improved fault tolerance: non-blocking mode, new
init function with config, and ncclCommFinalize function.
Reintroduce collnet+chain algorithm, alongside collnet+direct.
Add LL protocol for intra-node P2P (on by default) and network
communication (off by default).
Use network instead of shared memory when performance is better.
Fix: wait for CUDA graph destroy before destroying comm with linked
graph resources.
Remove aggressive polling during enqueue.
Fix DMABUF fallback on MOFED 5.4 and earlier.


[ROCm/rccl commit: c4e2aa6c79]
2022-08-18 02:53:17 -07:00
akolliasAMD 6fb5c5d5e3 minor latency tuning (#591)
* minor tuning for tree ll

[ROCm/rccl commit: 4cecdc9be5]
2022-08-03 15:07:44 -06:00
Sylvain Jeaugey 91154e8df9 2.13.4-1
Optimize CUDA graph launch; avoid launching a CPU callback for
intra-node operations.
Simplify kernel common code to improve the latency of send/recv
operations.
Strengthen CUDA streams semantics.
Change NET API to v6, to add dmabuf support.
Add ncclGetLastError() function.
Add ncclRemoteError code and use it for remote network errors.
Support the use of a different NCCL_NET parameter per communicator.
Add support for SHM and P2P transfers using cudaMemcpy.


[ROCm/rccl commit: 19ab67d172]
2022-07-11 08:10:34 -07:00
Wenkai Du 2f4aea93e0 Fix GPU to NIC mapping in tree (#573)
* Fix GPU to NIC mapping in tree

* Update tuning table

[ROCm/rccl commit: 00af1f64e9]
2022-07-03 20:52:52 -07:00
Wenkai Du 11a6cdd52f Fix P2P scheduling (#560)
[ROCm/rccl commit: 5cb2aca3d9]
2022-06-06 13:32:28 -07:00
Aristotelis 0b55e01ef3 Merge remote-tracking branch 'ncclRepo/master' into develop
[ROCm/rccl commit: e0864e7093]
2022-06-02 15:27:24 +00:00
akolliasAMD 22dc8bd246 Added creation of new tree and added switch for using treesplit for specific cases (#551)
[ROCm/rccl commit: 98f0809a39]
2022-05-25 18:55:14 -04:00
Sylvain Jeaugey 1c5734046d 2.12.12-1
Improve allreduce performance when we have more than one network interface per
GPU and we need to use PXN to close rings.
Add support for PCI Gen5 on 5.4 kernels.
Fix crash when setting NCCL_SET_THREAD_NAME.
Fix random crash in init due to uninitialized struct.
Fix hang on cubemesh topologies.
Add P2P_DIRECT_DISABLE parameter to disable direct access to pointers within a
process.


[ROCm/rccl commit: 7aa1c46fd5]
2022-05-13 00:26:57 -07:00
Wenkai Du 347ea354c2 Update tuning parameters
[ROCm/rccl commit: 83fd4f70e7]
2022-04-18 16:04:04 -07:00
Wenkai Du 67e7e6507e Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: d28e1cb44f]
2022-04-18 11:15:25 -07:00
Sylvain Jeaugey 27130280b2 2.12.10-1
Fix bug with CollNet
Fix bug with zero-bytes send/recv operations
Fix NCCL_PARAM implementation to avoid taking a lock on every call
Fix bug when setting NCCL_IB_QPS_PER_CONNECTION to more than one.
Improve error reporting for network errors.


[ROCm/rccl commit: 353e8ba446]
2022-03-30 02:27:01 -07:00
Wenkai Du 828f3d11a0 Update tuning parameters (#518)
* Update tuning parameters

* Respect user algo and topo selections

[ROCm/rccl commit: 7cbbca4da1]
2022-03-29 08:15:37 -07:00
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: 3c223c105a]
2022-03-02 20:48:56 +01:00
Ziyue Yang dfa9b9e958 Add Pivot AllToAll algorithm for Rome model (#503)
* add a2a pivot interface

* remove debug info

* address comments

* fix bug

* remove custom script

* address comments

* fix bug

[ROCm/rccl commit: b569c0a1db]
2022-02-20 21:09:47 -08:00
Wenkai Du 5b697e40db Add another Rome model (#497)
[ROCm/rccl commit: 02096c9936]
2022-02-12 10:30:16 -08:00
Wenkai Du 123de1621f Use hipGraph instead of cudaGraph (#487)
[ROCm/rccl commit: 3669e12432]
2022-01-10 08:26:01 -08:00
Wenkai Du f8bd2d0cfa Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 3a919c1f49]
2021-11-11 14:22:12 -08:00
Wenkai Du fc5f692f71 Remove extra work element copy (#465)
[ROCm/rccl commit: e05de8fd26]
2021-11-09 13:52:03 -08:00
John Bachan 389213c7e8 Fix compilation failure in "src/enqueue.cc" on older GCC because of
missing `#include <cstring>`.


[ROCm/rccl commit: 30ca3fcacf]
2021-09-23 09:55:16 -07:00
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: e11238b302]
2021-09-08 16:06:23 -07:00
John Bachan 04553b802a Fix to https://github.com/NVIDIA/nccl/issues/560
ncclGroup's containing operations of mixed datatype, element, or collective
would induce crash.


[ROCm/rccl commit: 5f2f2f670f]
2021-08-31 15:50:05 -07:00
Wenkai Du 95b37d97cf Merge remote-tracking branch 'nccl/master' into 2.10.3
[ROCm/rccl commit: 21969d7f89]
2021-09-01 16:34:44 -07:00
Wenkai Du 4b89e98675 Merge remote-tracking branch 'nccl/master' into 2.10.3
[ROCm/rccl commit: bf2339f93e]
2021-07-30 16:23:14 -07:00
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: 7e51592129]
2021-07-08 14:30:14 -07:00
Wenkai Du f82f99f533 Select sendrecv path based on collective data size (#391)
* Select sendrecv path based on collective data size

* Add comments on packing and unpacking group field

* Toggling RCCL_P2P_NET_DISABLE in combined calls unit tests

[ROCm/rccl commit: 6dcae8a459]
2021-06-10 17:51:04 -07:00
Wenkai Du 5bebcb0015 Setup collectives threshold for enabling intranet (#387)
* Setup collectives threshold for enabling intranet

* Use separate operation counters for coll and p2p

[ROCm/rccl commit: b815a2800f]
2021-06-09 13:24:26 -07:00
Wenkai Du b154b532a2 Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: e3abf1c2ec]
2021-05-25 20:52:15 -07:00
Wenkai Du b27490d38d Use fixed segment size for sendrecv (#369)
[ROCm/rccl commit: 8cde34be51]
2021-05-19 08:25:26 -07:00
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: 3fec2fa5ee]
2021-05-12 11:09:31 -07:00
gilbertlee-amd 071150a1b4 Clique tuning upgrade (#352) (#19)
* Enabling clique for any XGMI-connected topology, adding tuning
* Updating CHANGELOG for clique tuning
* Re-working clique barrier system to work on multi-process / multi-gpu

[ROCm/rccl commit: e796b1645c]
2021-05-11 08:44:59 -06:00
Wenkai Du a76bebf8b6 Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: a4ea1fed5b]
2021-05-05 16:01:01 -07:00
Sylvain Jeaugey 20da390b96 2.9.6-1
Add support for CUDA graphs.
Fuse BCM Gen4 switches to avoid suboptimal performance on some platforms. Issue #439.
Fix bootstrap issue caused by connection reordering.
Fix CPU locking block.
Improve CollNet algorithm.
Improve performance on DGX A100 for communicators with only one GPU per node.


[ROCm/rccl commit: a46ea10583]
2021-04-12 16:00:46 -07:00
Wenkai Du 065bde98d8 collnet: support multiple NICs (#335)
[ROCm/rccl commit: d87dc7c2e8]
2021-03-25 20:59:32 -07:00
Wenkai Du 287ed0f18a Enable collnet in RCCL (#333)
* Enable CollNet and use different number of channels

* topo_expl: enable collnet

[ROCm/rccl commit: 1d6244b18d]
2021-03-19 12:58:13 -07:00
Wenkai Du b7253710ca Revert "Port alltoall[v]" (#325)
This reverts commit 2c49121171.

[ROCm/rccl commit: 8e180cf087]
2021-03-06 13:59:31 -08:00