Wenkai Du
818cdb16a8
Query XGMI links from xml and adjust gfx906 channel usage ( #410 )
2021-07-27 17:32:41 -07:00
Wenkai Du
135d47d125
topo_expl: fix build after switching to rocm-smi-lib ( #405 )
...
* topo_expl: fix build after switching to rocm-smi-lib
* Use minimal of 4 channels for gfx908
2021-07-27 08:30:08 -07:00
Wenkai Du
fa6d7e9a63
Fixes for NCCL_MAX_NCHANNELS and topo_expl ( #398 )
2021-06-22 08:41:49 -07:00
Wenkai Du
b815a2800f
Setup collectives threshold for enabling intranet ( #387 )
...
* Setup collectives threshold for enabling intranet
* Use separate operation counters for coll and p2p
2021-06-09 13:24:26 -07:00
Wenkai Du
a3a8c2d56b
Allow intranode use of network connection ( #383 )
...
* Allow intranode use of network connection
* Checking for graph for null pointer
2021-06-08 07:37:59 -07:00
Wenkai Du
13dc80ee14
topo_expl: update to 2.9.9
2021-05-26 09:24:34 -07:00
Wenkai Du
a4ea1fed5b
Merge remote-tracking branch 'nccl/master' into develop
2021-05-05 16:01:01 -07:00
Wenkai Du
e26ad2995e
Cleanup number of channels calculation ( #340 )
2021-04-05 17:51:56 -07:00
Wenkai Du
17491c918e
Fix incorrect net counting ( #339 )
...
* Fix incorrect net counting
* Add comments
2021-04-05 12:21:57 -07:00
Wenkai Du
d87dc7c2e8
collnet: support multiple NICs ( #335 )
2021-03-25 20:59:32 -07:00
Wenkai Du
1d6244b18d
Enable collnet in RCCL ( #333 )
...
* Enable CollNet and use different number of channels
* topo_expl: enable collnet
2021-03-19 12:58:13 -07:00
Wenkai Du
8e180cf087
Revert "Port alltoall[v]" ( #325 )
...
This reverts commit f4d5d3d620 .
2021-03-06 13:59:31 -08:00
Wenkai Du
c018edf0f2
Enable local sendrecv over network if GDR is available on all GPUs ( #324 )
2021-03-05 19:59:41 -08:00
Wenkai Du
95f178324c
Add support to another Rome model
2021-02-18 02:00:31 +00:00
Wenkai Du
2ddbe6646b
Improve collective trace
2021-01-14 19:28:01 -05:00
Wenkai Du
f4d5d3d620
Port alltoall[v]
2021-01-14 19:28:01 -05:00
Wenkai Du
d469947641
Merge remote-tracking branch 'nccl/master' into no-target-id
2021-01-14 19:27:53 -05:00
Wenkai Du
ae008fd2db
Rework Rome detection and add multiple network ports models ( #274 )
...
* Rework Rome detection and add multiple network ports models
* Remove unused opCount in p2p transport
2020-10-07 13:37:36 -07:00
Wenkai Du
b871ea3c0c
Add Alltoallv RCCL kernel implementation ( #269 )
...
* Add alltoallv API and implementation
* Extend Rome P2P channel limit to multinode and alltoall kernels
* topo_expl: fix compilation and sync up with main
* gtest: use RCCL alltoallv API
* Code review changes
2020-09-30 16:25:36 -07:00
Wenkai Du
c5cbece6d0
Increase minimal channels for gfx908 ( #259 )
2020-08-26 11:40:11 -07:00
Wenkai Du
ab787c767e
Change default channels duplication for chordal ring ( #233 )
2020-07-14 15:16:50 -07:00
Wenkai Du
a6be82f5ab
topo_expl: fix broken build ( #224 )
2020-06-30 11:11:23 -07:00
Wenkai Du
e80e29573c
Add gather, scatter and alltoall collectives
...
Introducing 3 new APIs:
ncclResult_t ncclGather(const void* sendbuff, void* recvbuff, size_t sendcount,
ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream);
ncclResult_t ncclScatter(const void* sendbuff, void* recvbuff,
size_t recvcount, ncclDataType_t datatype, int root, ncclComm_t comm,
hipStream_t stream);
ncclResult_t ncclAllToAll(const void* sendbuff, void* recvbuff, size_t count,
ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream);
Only out of place operation is supported.
Preprocessor symbol RCCL_GATHER_SCATTER=1 indicates API availibility.
By default the APIs launche RCCL kernel implementation, which can be disabled by
RCCL_ALLTOALL_KERNEL_DISABLE=1. Then the APIs use wrapper around ncclSend and ncclRecv.
2020-06-09 17:44:08 -07:00
Wenkai Du
71ec3e09df
tpol_expl: update to 2.7
2020-06-09 17:40:24 -07:00
Wenkai Du
6f54b23503
topo_expl: update to 2.6
2020-04-01 13:37:08 -07:00
Wenkai Du
55f8e2dec7
Add topology explorer
2020-02-19 14:42:06 -08:00