38 Commits

Author SHA1 Message Date
Sourav Chakraborty 046af13751 Fix build failure in rccl_prim_test (#1984)
Added missing header in rccl_prim_test

[ROCm/rccl commit: 5b345d105c]
2025-10-21 12:51:14 -05:00
corey-derochie-amd e95578ef4c 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>

[ROCm/rccl commit: 6505639cf4]
2025-03-20 09:34:53 -06:00
Wenkai Du 890fafc2f7 rccl_prim_test: increase max number of workgroups and test iterations (#1132)
[ROCm/rccl commit: e8c76fd806]
2024-04-03 11:29:21 -07:00
Wenkai Du 4eaf90f84c p2p-latency-tests: fix build by switching to gcnArchName (#1030)
* p2p-latency-tests: fix build by switching to gcnArchName

* rccl-prim-test: switch to gcnArchName

[ROCm/rccl commit: cfc04a8aef]
2024-01-04 13:36:48 -08:00
Wenkai Du d2a1142cf3 rccl-prim-test: calculate iterations' standard deviation (#803)
* rccl-prim-test: calculate iterations' standard deviation

* Add default ring configuration for gfx940

* Use hipDeviceMallocUncached on gfx94x

[ROCm/rccl commit: f41ea11444]
2023-07-13 11:05:50 -07:00
Wenkai Du a5feda50b9 rccl-prim-test: calculate throughput standard deviations (#802)
[ROCm/rccl commit: 43f13cd25a]
2023-07-12 10:04:40 -07: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
Wenkai Du 6f3d035396 rccl-prim-test: minor update (#718)
[ROCm/rccl commit: addbf4bd90]
2023-04-03 07:30:04 -07:00
akolliasAMD 2cd63dac42 added stream synch after hipMemset (#609)
[ROCm/rccl commit: 06bce9d0c9]
2022-08-30 16:18:37 -06:00
Wenkai Du 310d51056f rccl-prim-test: enable 8p1h and 16p1h test
[ROCm/rccl commit: 7558b5e2bf]
2021-09-08 11:51:26 -05:00
Wenkai Du 4f610a2239 Revert "rccl-prim-test: add all-to-all benchmark (#185)"
This reverts commit e3e1c6b29c.


[ROCm/rccl commit: b22d097524]
2021-09-07 16:41:46 -05:00
Wenkai Du 661b1351a3 Limit max channels for ring graph on single node Rome (#347)
* Limit max channels for ring graph on single node Rome
* Partially revert "Use non-temporal access for streaming data (#341)"

[ROCm/rccl commit: a79f74082e]
2021-04-14 10:14:54 -07:00
Wenkai Du 6b3389b790 Use non-temporal access for streaming data (#341)
* Use non-temporal access for streaming data

* Revert to ulong2 after fixing compiling issue

[ROCm/rccl commit: 9dfc2c183e]
2021-04-07 17:34:35 -07:00
Wenkai Du adff98765c Merge remote-tracking branch 'nccl/master' into no-target-id
[ROCm/rccl commit: d469947641]
2021-01-14 19:27:53 -05:00
MurtadhaAldallal f1373612b0 Update rccl_prim_test.cpp (#246)
Adding doublelocalcopy operation and freeing buffer memory at end.
DoubleLocalCopy Patch Added

[ROCm/rccl commit: 390c63cf0d]
2020-08-07 08:20:14 -07:00
Stanley Tsang 3dd54437fe Adding appropriate references in rccl-prim-test (#227)
Adding appropriate references to rccl-prim-test.

[ROCm/rccl commit: 9bd4c14603]
2020-07-06 10:15:03 -06:00
Wenkai Du ced9958094 rccl-prim-test: add flags when calling hipExtLaunchMultiKernelMultiDevice in hip-clang
[ROCm/rccl commit: f1058b6353]
2020-05-12 23:54:07 +00:00
Wenkai Du 2de0b24c30 rccl-prim-test: auto-detect rings in 4P and 8P configurations
[ROCm/rccl commit: ef7064ba9b]
2020-04-10 18:17:21 +00:00
Wenkai Du e3e1c6b29c rccl-prim-test: add all-to-all benchmark (#185)
For gfx908, support simple detection of ring topology.
Call ReduceOrCopyMulti directly from kernel.
Also simplify code by removing kernel start synchronization option
which has no effect on throughput measurements.

[ROCm/rccl commit: ebc823e603]
2020-03-16 10:00:54 -07:00
Stanley Tsang e5419407c4 Updating copyright notices for 2020.
[ROCm/rccl commit: 20fa04d9b6]
2020-01-29 15:28:08 -08:00
Wenkai Du 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
Wenkai Du 25b3175e82 rccl_prim_test: Generalize ring topology and duplications
Allow user specified ring topology from command line and duplicated
to requested number of workgroups:
./rccl_prim_test -w 12 -p copy -r "0 1 2 3|3 2 1 0|0 2 1 3|3 1 2 0|0 2 3 1|1 3 2 0"


[ROCm/rccl commit: 07bb6fce8f]
2019-11-11 15:42:24 -08:00
Wenkai Du 62042e47bc Correct RTC frequencies for profiling purpose
[ROCm/rccl commit: 8995047830]
2019-11-05 11:36:45 -08:00
Wenkai Du 91b906cf88 rccl-prim-test: use hipExtLaunchMultiKernelMultiDevice and minor cleanup
[ROCm/rccl commit: ab91cdd5c9]
2019-10-30 13:15:02 -07:00
rohit pathania 31670133a1 Read operation throughput
[ROCm/rccl commit: a270ee080e]
2019-09-03 14:58:40 +05:30
rohit pathania bc51b5bc28 display each workgroup ,links and directions with throughputs
[ROCm/rccl commit: e5b13d69e5]
2019-08-30 13:28:23 +05:30
rpathani c441f2ff9b Update rccl_prim_test.cpp
[ROCm/rccl commit: 40e30b5168]
2019-08-19 12:44:11 +05:30
rpathani eaa1cdb48c Merge branch 'master' into xgmi_bench
[ROCm/rccl commit: deea20d49c]
2019-08-16 10:56:56 +05:30
Wenkai Du 6047487815 RCCL 2.4 update
[ROCm/rccl commit: f11c8f60cd]
2019-08-14 10:42:35 -07:00
rohit pathania 2dbcb62caf Modified the code to use RTC clock frequency based on gpu gcn id
[ROCm/rccl commit: 65e2f5d87b]
2019-08-14 12:55:12 +05:30
rohit pathania 042261445d Merge branch 'xgmi_bench' of https://github.com/rpathani/rccl into xgmi_bench
# Conflicts:
#	tools/rccl-prim-test/rccl_prim_test.cpp


[ROCm/rccl commit: 0f74929dab]
2019-08-13 11:36:56 +05:30
rohit pathania 86f6d95b06 Adding linkinfo and srcGPU to destGPU info
[ROCm/rccl commit: 3bbf924ff8]
2019-08-13 11:28:50 +05:30
rohit pathania 95162665c7 Adding linkinfo and srcGPU to destGPU info
[ROCm/rccl commit: 5a2f74b8d0]
2019-08-09 12:44:06 +05:30
Wenkai Du b7322c800a Refactor primitive test to support multiple GPUs in rings (#94)
* Refactor primitive test to support multiple GPUs in rings

* Make GPUs sync before transfer optional

* Use same ring format as RCCL

* Extend to 8 GPUs and report errors if there is no P2P access

* Control GPUs sync before ops from command line with "-s" option

* Change buffer size through command line option "-n"

Rename iterations command line option to "-i"


[ROCm/rccl commit: 70804da15b]
2019-07-05 14:29:20 -07:00
Wenkai Du 90b7a02944 Match primitives unroll counts with latest RCCL (#91)
[ROCm/rccl commit: e6a0da444f]
2019-06-26 15:09:13 -07:00
Wenkai Du 87d5441552 Calculate and print kernel throughput (#78)
* rccl-prim-test: print GPU info and set iterations

* Calculate and print kernel throughput


[ROCm/rccl commit: ee14676064]
2019-06-07 10:39:30 -07:00
Wenkai Du dcb2801f25 rccl-prim-test: print GPU info and set iterations (#77)
[ROCm/rccl commit: 42b488507d]
2019-06-05 15:16:33 -07:00
Wenkai Du 0ed10b1e4d Add RCCL primitive testing (#70)
[ROCm/rccl commit: 1bb6d2104c]
2019-05-23 16:52:17 -06:00