Commit Graph

510 Commits

Author SHA1 Message Date
Nusrat Islam 7ea099612f device: use unroll factor based on platforms
[ROCm/rccl commit: 3a741787bf]
2023-06-14 13:36:15 -05:00
Bertan Dogancay 1ae071944e improve compilation time and create timetrace plot (#773)
* improve compilation time and create time-trace plot

* set default value for nproc

[ROCm/rccl commit: f35777e9b0]
2023-06-14 09:17:51 -06:00
akolliasAMD 8438fd9e42 Wall clock update and npkit trace script Update (#771)
* changed builtin clock to wall_clock64
* updated npkit_Trace_generator to the new version of npkit

[ROCm/rccl commit: 9cdac774ea]
2023-06-07 17:47:10 -06:00
Cory Bloor 8df2b752cd Fix build on additional architectures (#740)
* Fix build on additional architectures

Instead of directly wrapping a platform-specific operation with a
preprocessor check against a gfx macro, it can be more flexible to
check a macro that can be overriden by the user. The gfx macro can then
just provide the default value for the macro, resulting in the same
default behaviour as if the gfx macro was checked directly but with
more control at build-time.

For example, to build rccl without using buffer_wbinvl1_vol on
gfx902, but still use the default on other archs, a user could
export CXXFLAGS='-Xarch_gfx902 -DRCCL_USE_WBINVL1_VOL=1' before
configuring the build. This flexibility isn't always necessary, but
it's nicer to have it and not need it than to need it and not have it.

* Define WARP_SIZE using warpSize builtin

[ROCm/rccl commit: b1a65afd58]
2023-06-06 16:45:50 -06: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
Bertan Dogancay a32eeae7b5 add DMA_BUF support (#763)
* add DMA_BUF support

* remove unused libraries in src/init.cc

* change NCCL_ALL to NCCL_INIT

* remove extra pointer functions in transport/net.cc

[ROCm/rccl commit: d52b6c0d24]
2023-06-01 12:46:42 -06:00
Wenkai Du 61e30182b2 Rework barrier and event code (#761)
* Rework barrier and event code

* Switch to inline asm

[ROCm/rccl commit: 5a38ff192b]
2023-05-31 13:36:51 -07:00
Nusrat Islam 325ff1dc11 device: change unroll factor
The default value of unroll factor is 2. Changing the unroll
factor to 4 provides better performance for most of the collectives.


[ROCm/rccl commit: 4d1cfb17c8]
2023-05-25 15:42:35 -05:00
Ziyue Yang a7557cf7b0 revert npkit (#748)
[ROCm/rccl commit: 7d6e7bcd7d]
2023-05-24 07:41:05 -07:00
Ziyue Yang 4430e4448f Limit MSCCL reduce unrolling to pow-2 cases to shrink kernel size (#746)
[ROCm/rccl commit: ed252c30f4]
2023-05-19 11:46:36 -07:00
Ziyue Yang 73c6d51454 fix min, max and avg (#745)
[ROCm/rccl commit: 11676267b5]
2023-05-18 11:02:59 -07:00
Wen-Heng (Jack) Chung 1bafe9c240 Merge pull request #742 from whchung/skip_done_event_msccl
Allow skipping doneEvent inside MSCCL.

[ROCm/rccl commit: eba4e9e100]
2023-05-18 10:17:20 -05:00
Wenkai Du f119977daa Fix merge error (#744)
[ROCm/rccl commit: 403cda6322]
2023-05-18 08:09:27 -07:00
Wen-Heng (Jack) Chung 932822fc46 Address review feedbacks and make the flag be disabled by default.
[ROCm/rccl commit: ca4a1dfd67]
2023-05-17 17:50:25 +00:00
Wen-Heng (Jack) Chung a7fbe68535 Skip doneEvent inside MSCCL by default.
Added a RCCL_MSCCL_ENABLE_DONE_EVENT env var, set it be 0 by default.

The env var is to control whether to use doneEvent when invoking MSCCL
kernels.

Skipping doneEvent would cause the firmware to skip L2 cache flush,
resulting in overall performance improvement.


[ROCm/rccl commit: 12dba425de]
2023-05-17 16:49:42 +00:00
Wenkai Du f7184aef3f Revert "Ensure memory copy integrity during transport setup (#731)" (#741)
* Revert "Ensure memory copy integrity during transport setup (#731)"

This reverts commit 0c0c35927a.

Add stream synchronization in ncclStrongStreamRelease.

* Use event record and wait

[ROCm/rccl commit: 4ca7742c61]
2023-05-16 10:34:47 -07:00
Wenkai Du 3c7cf5b110 Skip checking of some settings in Cray OS (#739)
[ROCm/rccl commit: 8bb3340fcb]
2023-05-09 07:59:56 -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 0c0c35927a Ensure memory copy integrity during transport setup (#731)
[ROCm/rccl commit: 36e453c61e]
2023-04-25 14:41:43 -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
Kaiming Ouyang 5eba5f178d Add a comment to shutdown() in ncclSocketClose
[ROCm/rccl commit: 006b6bc7dc]
2023-04-13 09:13:44 -07:00
Kaiming Ouyang e35b05a872 Shutdown socket before close in ncclSocketClose()
[ROCm/rccl commit: 367e9b61c3]
2023-04-13 09:11:52 -07:00
Ziyue Yang 29c3da724c MSCCL: Fix memcpy bug (#721)
[ROCm/rccl commit: 7289c05146]
2023-04-11 14:46:53 -07:00
Ziyue Yang aa605b0989 fix msccl stream usage (#717)
[ROCm/rccl commit: c8e33b1232]
2023-03-24 10:59:36 -07:00
Wenkai Du 5d9c3b0277 Fix unit test HIP graph error (#712)
[ROCm/rccl commit: b02fd04165]
2023-03-20 15:34:09 -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
Wenkai Du 62f5e6a82f Warn user on incorrect system settings (#696)
* Warn user on incorrect system settings

* Fix typo

* Add possible impact

* Ignore iommu settings in VM

[ROCm/rccl commit: 79a2031951]
2023-03-06 08:17:06 -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 8aaf6f8e0d Merge pull request #685 from ROCmSoftwarePlatform/2.16.5
Sync up to NCCL 2.16.5

[ROCm/rccl commit: d601c4909c]
2023-02-22 10:29:02 -08:00
Wenkai Du 1f42d485f8 Fix P2P scheduling (#690)
[ROCm/rccl commit: 86e7b71234]
2023-02-21 07:49:54 -08:00
Wenkai Du 393d0ba7f8 Add back __syncthreads() in barrier and adjust stack size (#688)
[ROCm/rccl commit: 1c166046a2]
2023-02-18 08:50:31 -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 0c7a94c462 Merge remote-tracking branch 'nccl/master' into HEAD
[ROCm/rccl commit: aee7b42bb8]
2023-02-14 17:14:13 -08: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 8a0d254c69 Add HIP event optimization and remove special code for gfx90a
[ROCm/rccl commit: 39534e8724]
2023-02-10 16:46:01 +00: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 d95a71b2ef Switch to hipLaunchHostFunc for HIP graph (#667)
[ROCm/rccl commit: 2288e9ae80]
2022-12-15 10:16:46 -08: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
Wenkai Du bb27f5c8ba Fix typo in detecting Intel platforms (#661)
[ROCm/rccl commit: b953544a59]
2022-12-07 13:36:11 -08:00
akolliasAMD 90de5e3a84 decreased warp size for gfx110x (#655)
[ROCm/rccl commit: eca623df07]
2022-12-01 12:19:21 -07:00
Wenkai Du 416b754b31 Query DMABuf support through HSA runtime API (#654)
[ROCm/rccl commit: fb9938cffa]
2022-11-30 08:53:03 -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 cb382a0a6e Fix incorrect rocm-smi ID conversion (#648)
[ROCm/rccl commit: 57764f8152]
2022-11-21 19:44:39 -08:00
Wenkai Du b72619495d Fix collective trace timestamp format (#647)
[ROCm/rccl commit: 9cb72a3d0f]
2022-11-21 08:11:12 -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