593d99d9a9
Check fine grained memory before enabling RDMA
...
Adding back the check which was lost from 2.5 merge.
[ROCm/rccl commit: fb59328a7b
2020-03-02 11:18:27 -08:00
2a66deb694
Merge pull request #179 from wenkaidu/search
...
Use fraction of system maxWidth as steps for searching
[ROCm/rccl commit: 8b5bc8bca2
2020-02-28 11:05:46 -08:00
b750defc28
Merge remote-tracking branch 'remotes/nccl/master'
...
[ROCm/rccl commit: 8e73a2ad60
2020-02-27 12:53:03 -08:00
45a7541582
Revise PCI BW numbers on Rome
...
[ROCm/rccl commit: d2adc61bf6
2020-02-26 13:17:49 -08:00
b4be0ff3b8
Use fraction of system maxWidth as steps for searching
...
This reverts previous workaround of deducting only half of width
from paths.
[ROCm/rccl commit: 8391637613
2020-02-26 09:14:35 -08:00
5747c3cac1
Fix abort handling in LL primitives
...
[ROCm/rccl commit: 077c3cda74
2020-02-25 13:42:54 -08:00
d640f38d56
Fix system maxSpeed and maxWidth calculation
...
[ROCm/rccl commit: 9b80b3633f
2020-02-24 15:18:57 -08:00
93d448e2fe
Fix incorrect CR8 detection
...
Also change level of ring graph print to help debugging
[ROCm/rccl commit: f54dc58113
2020-02-21 10:09:49 -08:00
cf4bce4ad3
Merge pull request #172 from wenkaidu/topo_expl
...
Add topology explorer
[ROCm/rccl commit: 5b3856f2ed
2020-02-20 15:16:55 -08:00
00f421ccbd
Add topology explorer
...
[ROCm/rccl commit: 55f8e2dec7
2020-02-19 14:42:06 -08:00
6034c27655
Fix Allgather operations above 4G with multiple GPUs per process.
...
Fixes nccl-tests#37.
Direct offsets were still on 32 bits in the low-level primitives.
[ROCm/rccl commit: c38f174bd4
2020-02-12 11:11:55 -08:00
8432e8a921
Generate 8G6L chordal ring from reference
...
[ROCm/rccl commit: abcfbf1231
2020-02-11 22:01:12 +00:00
6b2d7de200
Add ring bandwidth correction factor
...
[ROCm/rccl commit: d1dae2721d
2020-01-30 09:52:27 -08:00
e5419407c4
Updating copyright notices for 2020.
...
[ROCm/rccl commit: 20fa04d9b6
2020-01-29 15:28:08 -08:00
622b49e80a
Split primitive class to smaller structures
...
[ROCm/rccl commit: 486fd436af
2020-01-29 15:27:23 -08:00
d2fbcfea02
Misc fixes and improvements for 2.5.6
...
1. Fix RCCL unit test
2. Add ROME detection and tuning
3. Change default P2P level
4. Fix search algorithm for XGMI
5. Remove explicit channel duplication with implicit by using half of link speed
6. Add collective trace support
7. Correct Intel Skylake CPU detection and bandwidth
8. Fix topo connect function
9. Disable GDR read and remove unreachable code
10. Disable LL128 kernels
11. Add tuning parameters
12. Use original clock64() implementation which returns RTC counter value
13. Print out timestamp of collective trace
14. Do not use struct ncclColl in kernel launch parameter
15. Fix abort handling and add tracing
17. Add __launch_bounds__ to kernel functions
18. Remove unused abortCount
19. Unset default MIN_NRINGS and MIN_NCHANNELS
20. Do not allocate shared memory when not using LL128 kernels
21. Correct time print out in tuning log
[ROCm/rccl commit: 1e55645d97
2020-01-29 15:27:05 -08:00
d91217b16f
[topology] remove NET links when trimming system
...
This fixes a memory leak.
[ROCm/rccl commit: 7a18fe0784
2020-01-07 13:29:57 -08:00
ff74ebdcea
Fix clang build ( #274 )
...
The attribute is called `optnone`, not `noopt`.
[ROCm/rccl commit: 3899f6e0f2
2019-12-09 09:31:13 -08:00
e5a17ee58d
Fix clang compilation
...
[ROCm/rccl commit: aa15dfb29c
2019-12-06 09:55:54 -08:00
4984d5ce0b
Fix clang build ( #271 )
...
Clang doesn't understand `optimize("O0")`. It has `noopt`, which GCC doesn't understand. Wrap the difference in a macro.
[ROCm/rccl commit: 8c564e9b57
2019-12-06 09:14:55 -08:00
d7d4175df0
Merge remote-tracking branch 'remotes/nccl/master' into rccl_2.5.6
...
[ROCm/rccl commit: 6648c81dc6
2019-12-03 15:42:04 -08:00
4cb52294d5
Disable direct buffers to reduce scratch memory size
...
[ROCm/rccl commit: a0be2b8812
2019-11-20 13:03:16 -08:00
71560fd67b
2.5.6-1 ( #255 )
...
Add LL128 Protocol.
Rewrite the topology detection and tree/ring creation (#179 ). Improve
tree performance by sending/receiving from different GPUs. Add
model-based tuning to switch between the different algorithms and
protocols.
Rework P2P/SHM detection in containers (#155 , #248 ).
Detect duplicated devices and return an error (#231 ).
Add tuning for GCP
[ROCm/rccl commit: 299c554dcc
2019-11-19 14:57:39 -08:00
1e182391ad
Add bfloat16 support in RCCL
...
Preprocessor symbol RCCL_BFLOAT16 is used as feature indicator
[ROCm/rccl commit: 5e109ed400
2019-11-18 13:45:53 -08:00
62042e47bc
Correct RTC frequencies for profiling purpose
...
[ROCm/rccl commit: 8995047830
2019-11-05 11:36:45 -08:00
41f6319b33
Check for fine grain support using memory allocation
...
[ROCm/rccl commit: 669f1951a4
2019-11-01 15:58:49 -07:00
e43e1f1b3d
additional check for fine grain support in p2pCanConnect ( #146 )
...
[ROCm/rccl commit: 5a502955c9
2019-10-31 08:58:38 -07:00
d3f399f619
Disable HDP flush for RDMA
...
[ROCm/rccl commit: 296176a4fd
2019-10-23 14:40:17 -07:00
21bc1ef493
Revert collective chunk and slice steps to avoid drop in throughput
...
[ROCm/rccl commit: df74d12946
2019-10-18 12:54:00 -07:00
7560929bd7
Reverting GenericOp bug workaround modifications to slice/chunk steps
...
[ROCm/rccl commit: 37603ae6cb
2019-10-11 09:20:10 -07:00
cf597ff257
Performing __threadfence_system() with only first thread
...
[ROCm/rccl commit: 1392dd2997
2019-10-11 09:16:19 -07:00
d257970ad1
Fix for GenericOp device primitive bug
...
[ROCm/rccl commit: 8ae1bce3bb
2019-10-10 22:39:45 -07:00
fbcdfd8348
Merge pull request #136 from wenkaidu/tree
...
Enable tree kernels in build
[ROCm/rccl commit: 062c798c86
2019-10-09 10:58:52 -07:00
f86ee41415
Enable tree kernels in build
...
Need to tune and specify NCCL_TREE_THRESHOLD to allow usage
[ROCm/rccl commit: 76976c9e2e
2019-10-08 23:20:11 +00:00
d8a06589c9
Tuning the inline and unroll to reduce the scratch usage
...
Summary:
1. remove the noinline attribute for AllReduceThreeKernel;
2. change AUTPUNROLL for tree functions to 1 or 2;
Combining 1 and 2 will reduce the scratch usage from 1256 to 952
[ROCm/rccl commit: eec319038e
2019-10-08 14:02:25 -07:00
57dcac6afe
Only generate kernels for sum and copy
...
[ROCm/rccl commit: 61ef1dcad5
2019-09-24 17:01:12 -07:00
dbb7500fd1
RDMA HDP flush fix
...
[ROCm/rccl commit: 86ce0a93b5
2019-09-06 16:35:55 +00:00
2847ad9576
Revert "Set RDMA default to off state"
...
This reverts commit 4afd6818ba3e6b326a19
2019-09-05 18:16:53 +00:00
daf2c4b200
Allocate opCount in pinned host memory for P2P transport
...
To avoid remote P2P read access when checking remote GPU's opCount
[ROCm/rccl commit: 8c975353ed
2019-08-29 10:22:09 -07:00
4afd6818ba
Set RDMA default to off state
...
[ROCm/rccl commit: 0f16ad966a
2019-08-26 10:59:33 -07:00
4df1defc3b
Merge pull request #125 from wenkaidu/fix_nvml_id
...
Assign unused nmvlDev to avoid random number
[ROCm/rccl commit: 6759660529
2019-08-19 09:08:13 -07:00
04cd446d89
Assign unused nmvlDev to avoid random number
...
[ROCm/rccl commit: 86efdfc3b5
2019-08-16 16:34:14 -07:00
60989a3fc9
Merge remote-tracking branch 'remotes/nccl/master' into HEAD
...
[ROCm/rccl commit: 7c38da0939
2019-08-16 16:13:34 -07:00
7396d5c3ba
Tune AUTOUNROLL for better performance
...
Also remove all unused UNROLL defines
[ROCm/rccl commit: 1faededc03
2019-08-16 10:34:53 -07:00
f4a240065f
Fix build with hip-clang.
...
- Add necessary function attribute for HIP programming model.
- Explicitly include hsa headers.
[ROCm/rccl commit: 9369f8d75d
2019-08-15 14:56:04 -04:00
d4862fa605
Tune LL threshold for VEGA
...
Also move abort check after SPINS_BEFORE_CHECK_ABORT as NCCL
[ROCm/rccl commit: 2223cccf15
2019-08-15 09:16:11 -07:00
93c44e96cb
Default to minimal 2 rings and improve LL loop
...
[ROCm/rccl commit: 4b77a16f3f
2019-08-14 14:12:56 -07:00
1feef99e7d
Remove duplicate line
...
[ROCm/rccl commit: 5782a8d857
2019-08-14 13:22:43 -07:00
6047487815
RCCL 2.4 update
...
[ROCm/rccl commit: f11c8f60cd
2019-08-14 10:42:35 -07:00
d57c0b0f92
Updated PR#196 to use a common hash function
...
[ROCm/rccl commit: fad079a8ae
2019-08-14 10:08:39 -07:00
Wenkai Du
Wenkai Du
Sylvain Jeaugey
Stanley Tsang
Luke Yeager
Christian Sigg
Jeff Daily
Gilbert Lee
Changpeng Fang
Michael LIAO
David Addison