254208e7dd
Fixed issue with attomicEXCH creating errors on multi node runs ( #587 )
2022-07-22 11:32:49 -06:00
451c287aa6
Removing redundant LOAD and STORE on primitives plus adding some atomics ( #585 )
2022-07-21 13:04:57 -06:00
6dd090917a
Merge pull request #580 from hubertlu-tw/develop
...
Enhancement of RCCL logging information for topology-aware utilities
2022-07-15 15:16:37 -07:00
2b1d5d3bc1
intercept SIGUSR2 in RCCL
...
add support for intercepting SIGUSR2 in RCCL. This signal will
not terminate the execution of the application, but print the stacktrace
of the process that the signal was sent to instead.
2022-07-15 16:28:46 +00:00
77c2bef952
Remove redundant LOAD/STORE usage in LL initialization
2022-07-14 00:40:36 +00:00
a1842df858
Enhancement of RCCL logging information for topology-aware utilities
2022-07-11 19:01:10 +00:00
19ab67d172
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.
2022-07-11 08:10:34 -07:00
8c3c8b78c0
Skip HDP cache flush for gfx90a ( #578 )
...
* Skip HDP cache flush for gfx90a
* Remove extra debug print
2022-07-08 10:13:32 -07:00
aa0d7ca882
Add more constraints to enable GDR ( #579 )
...
* Add more constraints to enable GDR
* Revert deleted line
2022-07-08 09:52:27 -07:00
80f53cc171
Reduce AlltoAll port usage in send/recv proxy ( #577 )
...
* Reduce AlltoAll port usage when connecting proxy
Reuse socket ports when connecting proxies in AlltoAll.
Existing port usage in AlltoAll is O(n) for recv and O(n) for send,
reusing socket ports in server or client side will make one of them
O(1), reusing both will reduce the total port usage to O(1) and enables
AlltoAll in >64 MI200 nodes.
* Update changelog accordingly
Update changelog accordingly.
2022-07-07 16:15:52 -07:00
b250c01cbe
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
2022-07-06 07:58:41 -07:00
00af1f64e9
Fix GPU to NIC mapping in tree ( #573 )
...
* Fix GPU to NIC mapping in tree
* Update tuning table
2022-07-03 20:52:52 -07:00
a89a9966aa
Adding git hash info to version output line ( #572 )
2022-06-28 16:42:51 -06:00
9a285b5e1d
Do not set NET GDR level automatically ( #571 )
2022-06-23 16:28:28 -07:00
c3bb9e70d0
Use different atomics to check flags in kernel ( #568 )
2022-06-23 09:16:41 -07:00
6e93fafdc3
Add Feature - Add NPKit Support in RCCL ( #564 )
...
* apply npkit
* fix bug
* add npkit in readme
2022-06-20 14:30:19 -07:00
f274c865c1
Change default nchannels per peer ( #563 )
2022-06-13 06:39:05 -07:00
0336ffdf70
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.
2022-06-10 14:23:12 +00:00
5cb2aca3d9
Fix P2P scheduling ( #560 )
2022-06-06 13:32:28 -07:00
7a6c6927ae
Enable timing profile option ( #558 )
2022-06-03 07:05:13 -07:00
e0864e7093
Merge remote-tracking branch 'ncclRepo/master' into develop
2022-06-02 15:27:24 +00:00
eef812bed7
Revert chunksteps changes ( #555 )
2022-05-31 14:45:51 -07:00
ef499c4810
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
2022-05-31 11:31:30 -07:00
a0a686e74c
code cleanup ( #554 )
2022-05-31 09:59:36 -04:00
c5b77121f0
Update Rome model ( #552 )
2022-05-26 09:59:23 -07:00
98f0809a39
Added creation of new tree and added switch for using treesplit for specific cases ( #551 )
2022-05-25 18:55:14 -04:00
700b473211
Moving opt-in custom signal handler from UnitTests into RCCL ( #550 )
...
* Enable via RCCL_ENABLE_SIGNALHANDLER=1
2022-05-20 09:56:38 -06:00
6707a270b1
Add switch for pivot alltoall kernel ( #549 )
2022-05-17 18:14:04 -07:00
283dc86a73
Refine and add new Rome models ( #548 )
2022-05-17 08:23:59 -07:00
7aa1c46fd5
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.
2022-05-13 00:26:57 -07:00
c9919e0e35
Improve LL performance ( #546 )
...
* Improve LL performance
* Add split barriers for LL
2022-05-10 13:32:10 -07:00
83fd4f70e7
Update tuning parameters
2022-04-18 16:04:04 -07:00
d28e1cb44f
Merge remote-tracking branch 'nccl/master' into develop
2022-04-18 11:15:25 -07:00
fd2f1b3b88
Fix random segfault ( #537 )
2022-04-15 14:32:11 -07:00
2151c79d14
Add new Rome model ( #536 )
2022-04-13 11:45:40 -07:00
ba4c165bf3
Add new Rome model ( #535 )
2022-04-12 13:27:32 -07:00
9bfc1c6e35
Update Makefile to install static library.
...
Make sure make install also installs the static library.
Fixes #662
2022-04-08 14:00:43 +02:00
15b572751e
Increase chunk steps of broadcast and reduce ( #528 )
2022-04-07 13:34:04 -07:00
b2ffcf6d89
Doxygen fix for ncclRecv ( #527 )
...
Changed the Doxygen command for ncclRecv and pncclRecv, to be consistent with other APIs.
2022-04-05 14:07:56 -07:00
5cc0a405c0
Add tuning model ( #523 )
2022-04-04 10:19:57 -07:00
bbe780ca6c
Support multiple tuning tables ( #522 )
...
* Support multiple tuning tables
* [UnitTests] Skip managed memory testing
2022-03-31 17:09:21 -07:00
8133784b32
Merge remote-tracking branch 'origin/master'
2022-03-30 02:29:05 -07:00
353e8ba446
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.
2022-03-30 02:27:01 -07:00
2247152a8e
Fix merging error
2022-03-30 02:14:32 -07:00
2dfd83752c
Merge branch 'master' into truncated_msg_warning
2022-03-30 10:58:05 +02:00
1382a87306
Display host name instead of numeric IP when referring to a peer
...
For easier interpretation of debug messages like "connection closed by
peer", "peer message truncated" and "peer collective mismatch"
2022-03-30 10:47:10 +02:00
b895abcdb8
Fix typo in net_ib.cc
2022-03-30 10:45:01 +02:00
1c7c014ceb
Remove unnecessary newline in plugin logging
...
Signed-off-by: Felix Abecassis <fabecassis@nvidia.com >
2022-03-30 10:44:49 +02:00
7cbbca4da1
Update tuning parameters ( #518 )
...
* Update tuning parameters
* Respect user algo and topo selections
2022-03-29 08:15:37 -07:00
cd17cf6dce
Update Rome model matching and add new models ( #516 )
...
* Update Rome model matching and add new models
* Add missing file
* Models update
2022-03-21 10:54:40 -07:00
akolliasAMD
Hubert Lu
Edgar Gabriel
Ziyue Yang
hubertlu-tw
Sylvain Jeaugey
Wenkai Du
Yifan Xiong
gilbertlee-amd
Aristotelis
Wenkai Du
Colin Smith
Ke Wen
Christopher Hesse
Felix Abecassis