Gráfico de commits
277 Commits
Autor SHA1 Mensaje Fecha
Bertan Dogancay d411d52b19 Disable Colltrace for --fast option (#778)
* Disable Colltrace for --fast option

* Limit nprocs for CI

[ROCm/rccl commit: 0c77c66221]
2023-06-21 14:16:09 -06:00
Sylvain Jeaugey 2dc2c86ec1 2.18.3-1
Fix data corruption with Tree/LL128 on systems with 1GPU:1NIC.
Fix hang with Collnet on bfloat16 on systems with less than one NIC
per GPU.
Fix long initialization time.
Fix data corruption with Collnet when mixing multi-process and
multi-GPU per process.
Fix crash when shared memory creation fails.
Fix Avg operation with Collnet/Chain.
Fix performance of alltoall at scale with more than one NIC per GPU.
Fix performance for DGX H800.
Fix race condition in connection progress causing a crash.
Fix network flush with Collnet.
Fix performance of aggregated allGather/reduceScatter operations.
Fix PXN operation when CUDA_VISIBLE_DEVICES is set.
Fix NVTX3 compilation issues on Debian 10.


[ROCm/rccl commit: ea38312273]
2023-06-14 01:29:17 -07:00
Wenkai Du 90cbef7042 Add NCCL_NCHANNELS_PER_PEER override (#767)
Also fix topol_expl build issue

[ROCm/rccl commit: 3af90902c8]
2023-06-06 08:41:38 -07:00
Wenkai Du 18562abdb2 Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 53a1f91857]
2023-04-25 15:38:32 -07:00
Sylvain Jeaugey 902ff02645 2.18.1-1
Add support for IB SHARP to NVLS (NVLink SHARP algorithm).
Add NVLS+Tree algorithm.
Add support for memory management using cuMem* functions.
Use all NICs for Send/Receive operations on systems with more than
one NIC per GPU (#804).
Add ncclCommSplit primitive, with resource sharing option in config.
Fix alltoallv hang (#788)
Increase number of channels on H100 when we're not limited by NVLink.
Improve error reporting in case of IB failure, printing local and
remote ID (#779).
Add build option to allow compilation against RDMA includes instead
of dynamically loading IB verbs symbols (#802).
Fix context creation for progress thread (#803).
NET/IB: add option to use multiple QPs in round-robin mode.
Fix tree performance issue when NVB is disabled on HCM topologies.


[ROCm/rccl commit: d97a32fac8]
2023-04-18 03:58:25 -07:00
Ziyue Yang f7f669e7f0 MSCCL: Improve executor and integrate scheduler (#694)
* MSCCL: improve executor and add scheduler for testing

* Use external scheduler

* Fix cmake error

* Address comments

* Fix thread safe issue

* Make MSCCL lifecycle APIs thread safe

* Make MSCCL internal scheduler aware of topology hint

* Revise error message

[ROCm/rccl commit: e3b2342f39]
2023-03-14 14:34:25 -07:00
Wenkai Du 216c83c39d Fix XGMI detection (#699)
* Fix XGMI detection

* Increase stack size

* Temporarily disable signal hangler in CI

[Process: 17281] Inside handler function signal: Segmentation fault (11)
BFD: DWARF error: section .debug_info is larger than its filesize! (0x93ef57 vs 0x530ea0)
BFD: DWARF error: section .debug_info is larger than its filesize! (0x93ef57 vs 0x530ea0)

[ROCm/rccl commit: 22b81fbaae]
2023-03-08 14:08:07 -08:00
Sylvain Jeaugey 8dcf8e8720 2.17.1-1
Add new NVLS algorithm for allreduce using NVLink SHARP (intra-node only).
Add new config options: cgaClusterSize, minCTAs, maxCTAs, netName.
Enable LL128 when we use PXN to close rings.
NVTX3 includes update.
Fix crash when one CollNet (SHARP) rail fails to initialize.


[ROCm/rccl commit: 5d3ab08b69]
2023-03-01 00:39:04 -08:00
Wenkai Du 0c7a94c462 Merge remote-tracking branch 'nccl/master' into HEAD
[ROCm/rccl commit: aee7b42bb8]
2023-02-14 17:14:13 -08:00
Wenkai Du c76bc214c8 Merge remote-tracking branch 'nccl/master' into HEAD
[ROCm/rccl commit: e1cb45ff22]
2023-02-04 01:44:43 +00:00
Sylvain Jeaugey e6e8f2555c 2.16.5-1
Add support for 400Gbit NDR network adapters (CX7)
Handle EINTR in socket poll() function
Add NCCL_PROGRESS_APPENDOP_FREQ to control op append overhead
Resource cleanup fixes
Fix double free in case of init failure
Fix crash in ncclCommAbort
Revert AMD speed commit


[ROCm/rccl commit: f3d5166783]
2023-02-02 12:52:47 -08:00
Wenkai Du bb27f5c8ba Fix typo in detecting Intel platforms (#661)
[ROCm/rccl commit: b953544a59]
2022-12-07 13:36:11 -08:00
Sylvain Jeaugey 2ce8946622 2.16.2-1
Add support for CUDA 12.0, drop Kepler (sm_35).
Support for H100 features.
Make socket code more robust and protected. Solves #555.
Improve performance on large CUDA graphs, reducing dependencies.
Reduce inter-socket bandwidth on AMD CPUs to favor better paths.
Various fixes to ncclCommAbort.
Make service thread polling resistant to EINTR.
Compile with profiling API by default.
Extend NVTX instrumentation with call arguments.


[ROCm/rccl commit: 28189e2df8]
2022-11-30 02:31:59 -08:00
Wenkai Du 55e1008756 Adjust P2P channels on Intel platform (#653)
[ROCm/rccl commit: 9594bbee3b]
2022-11-29 13:57:10 -08:00
Wenkai Du ba18b77b07 Fix typo in previous hipify change (#645)
[ROCm/rccl commit: cf3c32a626]
2022-11-15 11:51:47 -08:00
Wenkai Du c15a10a9d2 Move hipify to cmake stage
Add minimal ROCm/HIP version requirements for Graph support


[ROCm/rccl commit: 562dd87036]
2022-11-14 18:10:45 +00:00
Wenkai Du ffecb74b1e Update tuning table and fix topo_expl
[ROCm/rccl commit: 94ad7f6f51]
2022-11-07 18:24:24 +00:00
Wenkai Du 4c9c1d41ee Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 9a077e6947]
2022-11-03 21:17:42 +00: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
Wenkai Du 76414d3230 Support P2P with invisible devices (#636)
* Support P2P with invisible devices

* Update copyright year

[ROCm/rccl commit: 9ddf0e0649]
2022-10-17 10:24:59 -07: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 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
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
Wenkai Du 6c3f1366e8 Add XGMI sys type and clean up detection code (#597)
[ROCm/rccl commit: f5c0b243a8]
2022-08-12 09:52:29 -07:00
Wenkai Du f70830d629 Revert "Use nontemporal in slow path and add XGMI sys type (#575)" (#590)
This reverts commit e04bba619a.

[ROCm/rccl commit: 9089e68a99]
2022-08-02 09:31:53 -07: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 659cd52d5c Add more constraints to enable GDR (#579)
* Add more constraints to enable GDR

* Revert deleted line

[ROCm/rccl commit: aa0d7ca882]
2022-07-08 09:52:27 -07:00
Wenkai Du e04bba619a Use nontemporal in slow path and add XGMI sys type (#575)
* Use nontemporal in slow path and add XGMI sys type

* Clean up XGMI detection

[ROCm/rccl commit: b250c01cbe]
2022-07-06 07:58:41 -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 915a9d3934 Do not set NET GDR level automatically (#571)
[ROCm/rccl commit: 9a285b5e1d]
2022-06-23 16:28:28 -07:00
Wenkai Du 0fb000932f Change default nchannels per peer (#563)
[ROCm/rccl commit: f274c865c1]
2022-06-13 06:39:05 -07:00
Edgar 8953f5b5ca Introduce multi-rank support per device.
This is a single commit of the source code changes required to
introduce support for multiple ranks per device.
A new interface (ncclCommRankInitMulti) has to be used to make use of
this new feature.


[ROCm/rccl commit: 0336ffdf70]
2022-06-10 14:23:12 +00:00
Aristotelis 0b55e01ef3 Merge remote-tracking branch 'ncclRepo/master' into develop
[ROCm/rccl commit: e0864e7093]
2022-06-02 15:27:24 +00:00
Wenkai Du 5becf1669f Add another Rome model (#553)
* Add another Rome model

* Add option to force enable intranet on single node

* Limit p2p channels to number of ranks

* Refine p2p channels handling

[ROCm/rccl commit: ef499c4810]
2022-05-31 11:31:30 -07:00
akolliasAMD a03ab8e752 code cleanup (#554)
[ROCm/rccl commit: a0a686e74c]
2022-05-31 09:59:36 -04:00
Wenkai Du 2c125ce6ed Update Rome model (#552)
[ROCm/rccl commit: c5b77121f0]
2022-05-26 09:59:23 -07: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
Wenkai Du b30b8becea Refine and add new Rome models (#548)
[ROCm/rccl commit: 283dc86a73]
2022-05-17 08:23:59 -07: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
Wenkai Du 58b2c1ec9c Fix random segfault (#537)
[ROCm/rccl commit: fd2f1b3b88]
2022-04-15 14:32:11 -07:00
Wenkai Du 011447e4dc Add new Rome model (#536)
[ROCm/rccl commit: 2151c79d14]
2022-04-13 11:45:40 -07:00