f3f3336468
Fix topo explorer's compatibility with NCCL 2.24 ( #1671 )
...
* Fix build issues
* Fix failure to find path remote rank
2025-05-05 15:26:29 -04:00
acfac55516
[Graph] Try using P2P by default ( #1670 )
2025-05-02 11:54:30 -04:00
329e13efff
Revert "[SRC] Enable unroll=1 for gfx950 ( #1602 )" ( #1667 )
...
* Revert "[SRC] Enable unroll=1 for gfx950 (#1602 )"
This reverts commit 307bc10781Nilesh.Negi@amd.com >
2025-04-30 23:33:08 -05:00
dc0c5f9153
minor fix for empty scope (group) ( #1666 )
2025-04-30 13:29:13 -04:00
a6bf9bfc9e
Merge remote-tracking branch 'nccl/master' into develop
2025-04-23 20:47:43 -07:00
82afb2bcfe
Expose production tuning table in topo_explorer using internal RCCL/NCCL logic ( #1628 )
...
* Internal RCCL/NCCL functionality exposed when RCCL_EXPOSE_STATIC is enabled
* Algo/protocol/max channels can be obtained with the new RCCL API
* Introduce rccl_static and rccl_static_inline macros to work around invisible functions in core source files like enqueue.cc
* Add usage example in topo-explorer tool
2025-04-23 15:44:56 -04:00
ac8ec4c08c
Fix NPKit for SendRecv ( #1651 )
2025-04-21 12:34:47 -04:00
9a55ff60a9
RCCL Replayer update ( #1603 )
...
RCCL recorder w/ suggested change and UT
2025-04-19 00:21:27 -04:00
52bfdf05dc
Address nested designator compiler warning issue ( #1633 )
2025-04-18 17:09:50 -04:00
f20c33effd
Fix MSCCLPP accuracy issue for allreduce7 ( #1634 )
...
* ext-src: fix a graph-mode bug in allreduce7
* change MSCCLPP threshold to 16MB
* ext-src: change message size threshold for allreduce7
* ext-src: address review comments
2025-04-18 08:54:32 -05:00
7a84c5dbb0
added copyright ( #1635 )
2025-04-14 09:46:18 -05:00
1786c0268b
Adding #include <dlfcn.h> in profiler.cc to pass build ( #1632 )
...
w/o the header, experience following errors:
```
error: use of undeclared identifier 'RTLD_NOW'
error: use of undeclared identifier 'RTLD_LOCAL'
error: use of undeclared identifier 'dlerror'
error: use of undeclared identifier 'dlsym'
error: use of undeclared identifier 'dlclose'
```
2025-04-10 08:48:18 -07:00
52654e2301
Add device synchronization before destroying proxy thread. ( #1631 )
...
This commit ensures that GPU finishes all kernel before destroying
communicator thread.
2025-04-10 10:44:16 -05:00
e40ff4f84a
all_reduce LL/LL128 and Ring/Tree multi-node tuning for MI300 ( #1627 )
...
* Enabling LL128 by default on MI300
* Add missing CUDACHECK
* Adjust BW correction factors to fix the Tree->Ring switching point
* Refactor and add ll128 AR logarithmic factor to tuning models
* Move RCCL tuning changes to a separate file
* Use enum for tunable indexing
* Use explicit indexing in tuning models to avoid mismatch issues
* Place rcclGetSizePerRank in a function
* Remove HIP ifdef for rccl-only call
---------
Co-authored-by: Mustafa Abduljabbar <mustafa.abduljabbar@amd.com >
2025-04-10 11:43:54 -04:00
4be06f04d8
Add AllGather LL128 multi-node tuning and include LL cutoff points in tuning models ( #1618 )
...
* Enable LL/LL128 cutoff points in tuning models
* Initializing ll/ll128 model cutoffs for MI300
* Use RCCL_LL_LIMITS_UNDEFINED
---------
Co-authored-by: PedramAlizadeh <pmohamma@amd.com >
2025-04-02 16:26:23 -04:00
b17338d164
Revert "[GRAPH] Increase default nChannels to 112 for gfx950 ( #1596 )" ( #1620 )
...
* Revert "[GRAPH] Increase default nChannels to 112 for gfx950 (#1596 )"
This reverts commit 1df73e209e
2025-03-28 17:57:06 -05:00
532f54c244
Merge pull request #1559 from BertanDogancay/2.23
...
[SYNC] 2.23.4-1
2025-03-28 17:06:56 -04:00
307bc10781
[SRC] Enable unroll=1 for gfx950 ( #1602 )
...
* [SRC] Enable unroll=1 for gfx950
* Fix typo from rebase in generate.py
* Support for unroll=1 and gfx90a when building for all GPU targets
---------
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com >
2025-03-27 18:21:35 -05:00
0b2062c560
Merge remote-tracking branch 'nccl/master' into develop
2025-03-27 12:53:04 -05:00
9dc23d9265
Disable mscclpp ( #1614 )
...
* disable mscclpp by default
2025-03-25 15:21:16 -05:00
626dc50ab5
Removing the experimental clique kernel files ( #1610 )
2025-03-20 18:10:01 -06:00
90ad586d94
Add fault injection of starting warps with random variations ( #1593 )
...
* Add fault injection of starting warps with random variations
This is done by inserting randomly delays after __syncthreads().
The feature can be turned off by FAULT_INJECTION=OFF in cmake.
* Remove manually introduced bug for demo purpose
* Use only one thread per warp for checking wall clock
2025-03-20 16:11:43 -07:00
6505639cf4
removed gfx940 and gfx941 ( #1606 )
...
* removed gfx940 and gfx941
* removed gfx940 and gfx941
* Update "gfx94" to "gfx942" in init.cc
* Updated remaining "gfx94" updates to "gfx942"
* Update filenames and variables from gfx940 to gfx942
---------
Co-authored-by: akolliasAMD <akollias@amd.com >
2025-03-20 09:34:53 -06:00
bd0092e8f1
GDRCOPY support: Off by default ( #1605 )
2025-03-18 08:17:01 -07:00
ccb0820743
Memory leak fix when numIBDevices = 0 ( #1429 )
...
* Initial commit for testing
* Fix memory leak in checkOptions
* Fix memory leak in checkOption
* x
* Delete cmake-3.28.2-linux-x86_64.sh
* gcn changes
* gcn memleak fixes
* gcn leak fix
* memory leak fixes for parseRome4P2H and ncclTopoAddGPU
* Keeping only necessary file for fixes
Deleting temporary scripts I created for debugging and testing
* changing to GCN_ARCH_NAME_LEN
* Added sanity check directory
* refactoring scripts
* Updated to sanity checks folder
* Initial fixes
* changes in tools
* pointing RCCL lib build to debug version
* Removed second pthread_detach
* Removing sanity checks
* Keeping only code changes
* addressing memory leaks in ncclIbinit
---------
Co-authored-by: Chao Chen <cchen104@amd.com >
2025-03-17 11:21:19 -05:00
f67b2cc908
Multi-node reduce_scatter improved auto-selection for LL and LL128 on gfx942 ( #1604 )
...
* Add reduce_scatter LL and LL128 thresholds
* Always honor user choice for protocol
2025-03-17 11:21:01 -04:00
245c2de909
Enable LL128 on gfx942 ( #1549 )
2025-03-16 15:10:05 -07:00
1df73e209e
[GRAPH] Increase default nChannels to 112 for gfx950 ( #1596 )
...
Signed-off-by: nileshnegi <Nilesh.Negi@amd.com >
2025-03-14 14:47:03 -07:00
4237caad69
Limit P2P channels per peer to not exceeding max channels ( #1594 )
...
* Limit P2P channels per peer to not exceeding max channels
* [UT] test single GPU cases for all collectives
* [UT] fix out of range root value
2025-03-11 09:32:09 -07:00
f6c6d451a9
mscclpp compatibility check for ubr ( #1573 )
2025-03-09 22:10:47 -04:00
ac823818aa
misc/msccl: force use of mscclpp ( #1581 )
2025-03-04 12:48:59 -06:00
d88cca3098
[Transport] Fix IntraNet ( #1582 )
2025-03-04 13:30:36 -05:00
f957c4fe22
NPKit: enable reduce scatter profiling ( #1580 )
2025-03-04 10:03:56 -08:00
23c0b7bd84
misc/msccl: Read graph capture status for every collective call ( #1576 )
...
* misc/msccl: read graphCaptureStatus for every collective call
* fix a bug in checking whether UBR is enabled in MSCCLPP
* cmake: Fix patch reversal order
* misc/msccl: add logging
2025-02-28 17:16:07 -06:00
60c1264d27
Improve RDMA flushing by write dummy payload with RO=0 ( #1570 )
...
* Improve RDMA flushing by write dummy payload with RO=0
* Rename env var for disabling this change to RCCL_GDR_FLUSH_GPU_MEM_NO_RELAXED_ORDERING
2025-02-27 16:20:32 -08:00
02b8e504f0
[BUILD] Fix generate.py for gfx950 ( #1575 )
...
* [BUILD] Fix generate.py for gfx950
Signed-off-by: nilenegi <Nilesh.Negi@amd.com >
* [BUILD] Cleaner way to check for gfx targets
Signed-off-by: nilenegi <Nilesh.Negi@amd.com >
---------
Signed-off-by: nilenegi <Nilesh.Negi@amd.com >
2025-02-26 18:43:37 -06:00
85eb1f16bc
Use bit reversal based mapping for multi-node ( #1572 )
2025-02-26 09:48:03 -05:00
f268553ee4
enable building rccl for gfx950 ( #1571 )
2025-02-25 16:13:48 -05:00
ddc5d58b93
Rail optimized trees ( #1540 )
...
* Allow disabling rail-optimized trees via RCCL_DISABLE_RAIL_TREES, Graphviz-friendly output via RCCL_OUTPUT_TREES
2025-02-20 15:18:29 -07:00
aedbc95735
reverted the syncLDS back to syncthreads ( #1554 )
2025-02-19 10:44:32 -07:00
baaa2ac64d
Insert barrier after loading work items to LDS ( #1551 )
2025-02-18 10:17:27 -08:00
32dc7ef47c
Enable GDRCopy only on gfx94x ( #1550 )
...
* Enable GDRCopy only on gfx94x
* Use cudaFree instead of hipFree
* Add warning if failed to get device property
* Remove extra return
2025-02-17 13:28:19 -08:00
0e5f4d0662
reverting the (Reduce NPKit latency overhead in MSCCL kernel) PR #893 ( #1525 )
2025-02-14 11:03:43 -05:00
824b81c034
Revert "replacing rccl_float8 with hip_fp8 and address compatibility issue (#…" ( #1545 )
...
This reverts commit d437d6e41c
2025-02-13 10:00:22 -07:00
d437d6e41c
replacing rccl_float8 with hip_fp8 and address compatibility issue ( #1538 )
...
* replacing rccl_float8 with hip_fp8 and address compatibility issue with gfx942
---------
Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com >
Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com >
2025-02-13 10:34:17 -06:00
ebf7e2305e
Print KL/CL/KE events for all warps ( #1544 )
...
* Print KL/CL/KE events for all warps
* Fix count off-by-one issue
* Fix opCount in KE and restore CPU thread option
* Simplify count calculation
2025-02-12 13:36:31 -08:00
f5b15f27a9
Move collective trace to HBM and fix log issue ( #1542 )
2025-02-11 11:40:14 -08:00
e8fb1335fd
Replace atomicAdd with _hip_atmoc_fetch_add in getting colltrace tail position ( #1539 )
2025-02-10 08:53:25 -08:00
3494f52d40
Adding AINIC Network Plugin check ( #1528 )
...
- Adding AINIC network plugin check to pass unused parameter to pass the channelId to the network plugin layer
2025-02-06 23:37:53 -06:00
a12bf32475
Reset barrier and make barrier_next thread local ( #1531 )
2025-02-05 09:06:48 -08:00
Mustafa Abduljabbar
Bertan Dogancay
Nilesh M Negi
Tim
BertanDogancay
Nusrat Islam
AbandiGa
Dingming Wu
Arm Patinyasakdikul
Pedram Alizadeh
isaki001
gilbertlee-amd
Wenkai Du
corey-derochie-amd
Avinash
akolliasAMD
mberenjk
Vijay Srinivasan