Commit graph

158 Commits

Autor SHA1 Nachricht Datum
Ziyue Yang a7557cf7b0 revert npkit (#748)
[ROCm/rccl commit: 7d6e7bcd7d]
2023-05-24 07:41:05 -07:00
Ziyue Yang 73c6d51454 fix min, max and avg (#745)
[ROCm/rccl commit: 11676267b5]
2023-05-18 11:02:59 -07:00
Wenkai Du a8f43b496c Remove references to NVLS functions
[ROCm/rccl commit: 897745a266]
2023-05-05 07:55:20 -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
Wenkai Du 1af52edd6a msccl: print stack and memory usage (#723)
* msccl: print stack and memory usage

* Update number of kernels calculation

[ROCm/rccl commit: 4b09ffba43]
2023-04-14 14:59:03 -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
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
Ziyue Yang 7c1290f995 NPKit: improve clock calibration and fix GPU clock API (#683)
* Improve clock calibration in NPKit

* Improve gfx macro

* Fix macro

[ROCm/rccl commit: f4bf47f325]
2023-02-17 12:26:57 -07:00
Wenkai Du 4fb1ebcf4b Remove workaround and use indirect function call (#684)
[ROCm/rccl commit: f7a456122c]
2023-02-14 13:59:48 -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
Ziyue Yang 431c7d2cb9 Add MSCCL Support (#658)
* Add MSCCL support

* Add alignment and message size checking

* Fix nRanks checking, in-place and out-of-place tests and group call handling

* Fix hipGraph unit test

* Change MSCCL init warning to INFO

* Revise license info

[ROCm/rccl commit: adafc0f759]
2022-12-12 15:51:04 -08:00
akolliasAMD 90de5e3a84 decreased warp size for gfx110x (#655)
[ROCm/rccl commit: eca623df07]
2022-12-01 12:19:21 -07: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 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 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
Ziyue Yang d54574b0eb apply npkit
[ROCm/rccl commit: 7d6bbc19d4]
2022-10-14 01:28:17 +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
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
Wenkai Du ba6e2db70d Misc fixes and disable binTree
[ROCm/rccl commit: a06e14e39b]
2022-09-14 00:26:19 +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
Gilbert Lee 1d24c476f4 Merge branch 'develop' into 2.13.4
[ROCm/rccl commit: 009e79623f]
2022-09-09 23:07:04 +00:00
gilbertlee-amd 35872115f8 Updating stream caching (#614)
- Adding non-captured hipStream for use in setup

[ROCm/rccl commit: dd56135a9a]
2022-09-09 16:30:15 -06: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
akolliasAMD 2cd63dac42 added stream synch after hipMemset (#609)
[ROCm/rccl commit: 06bce9d0c9]
2022-08-30 16:18:37 -06:00
Cosmic Fusion 1dcf1da5ca fix error: use of undeclared identifier 'free'
include stdlib.h to fix compilation error in rccl :

[39/58] Building CXX object CMakeFiles/rccl.dir/src/misc/signals.cc.o
FAILED: CMakeFiles/rccl.dir/src/misc/signals.cc.o 
/opt/rocm/bin/hipcc -DENABLE_COLLTRACE -DHAVE_BFD -DHAVE_CPLUS_DEMANGLE -DUSE_ROCM_SMI64CONFIG -D__HIP_PLATFORM_AMD__=1 -D__HIP_PLATFORM_HCC__=1 -Drccl_EXPORTS -I/home/cosmo/build/flgrwqa/build/include -I/home/cosmo/build/flgrwqa/build/include/rccl -I/home/cosmo/build/flgrwqa/rccl/src -I/home/cosmo/build/flgrwqa/rccl/src/include -I/home/cosmo/build/flgrwqa/rccl/src/collectives -I/home/cosmo/build/flgrwqa/rccl/src/collectives/device -I/opt/hsa/include -fPIC -fvisibility=hidden -fgpu-rdc -parallel-jobs=8 -Wno-format-nonliteral -x hip --offload-arch=gfx803 --offload-arch=gfx900:xnack- --offload-arch=gfx906:xnack- --offload-arch=gfx908:xnack- --offload-arch=gfx90a:xnack- --offload-arch=gfx90a:xnack+ --offload-arch=gfx1030 -std=c++14 -MD -MT CMakeFiles/rccl.dir/src/misc/signals.cc.o -MF CMakeFiles/rccl.dir/src/misc/signals.cc.o.d -o CMakeFiles/rccl.dir/src/misc/signals.cc.o -c /home/cosmo/build/flgrwqa/rccl/src/misc/signals.cc
In file included from /home/cosmo/build/flgrwqa/rccl/src/misc/signals.cc:8:
/home/cosmo/build/flgrwqa/rccl/src/include/BfdBacktrace.hpp:138:9: error: use of undeclared identifier 'free'
        free(file->syms);
        ^
/home/cosmo/build/flgrwqa/rccl/src/include/BfdBacktrace.hpp:155:5: error: use of undeclared identifier 'free'
    free(file->syms);
    ^

[ROCm/rccl commit: 080fc2d9d6]
2022-08-19 20:25:06 +03:00
Wenkai Du c2e9ada40b Repurpose profiling implementation to simple timestamps tracing (#600)
[ROCm/rccl commit: 14b8ff153f]
2022-08-18 15:34:46 -07: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 18d9fd1b8f Removing redundant LOAD and STORE on primitives plus adding some atomics (#585)
[ROCm/rccl commit: 451c287aa6]
2022-07-21 13:04:57 -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
Yifan Xiong bf15ad1d72 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.

[ROCm/rccl commit: 80f53cc171]
2022-07-07 16:15:52 -07:00
gilbertlee-amd cb5ae7224e Adding git hash info to version output line (#572)
[ROCm/rccl commit: a89a9966aa]
2022-06-28 16:42:51 -06:00
Ziyue Yang 2b418b5dee Add Feature - Add NPKit Support in RCCL (#564)
* apply npkit

* fix bug

* add npkit in readme

[ROCm/rccl commit: 6e93fafdc3]
2022-06-20 14:30:19 -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
Wenkai Du f2dbc77afe Enable timing profile option (#558)
[ROCm/rccl commit: 7a6c6927ae]
2022-06-03 07:05:13 -07: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 1e36b432f1 Revert chunksteps changes (#555)
[ROCm/rccl commit: eef812bed7]
2022-05-31 14:45:51 -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
gilbertlee-amd a2a4888497 Moving opt-in custom signal handler from UnitTests into RCCL (#550)
* Enable via RCCL_ENABLE_SIGNALHANDLER=1

[ROCm/rccl commit: 700b473211]
2022-05-20 09:56:38 -06: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 67e7e6507e Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: d28e1cb44f]
2022-04-18 11:15:25 -07:00
Wenkai Du 5ccdd9f5e1 Increase chunk steps of broadcast and reduce (#528)
[ROCm/rccl commit: 15b572751e]
2022-04-07 13:34:04 -07:00
Wenkai Du 3332cdff07 Support multiple tuning tables (#522)
* Support multiple tuning tables

* [UnitTests] Skip managed memory testing

[ROCm/rccl commit: bbe780ca6c]
2022-03-31 17:09:21 -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
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