AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge remote-tracking branch 'nccl/master' into 2.8.3

Browse Source 
[ROCm/rccl commit: c985358e11]
This commit is contained in:
Wenkai Du
2021-02-15 19:29:45 +00:00
parent ea2e665834 fc7bdb38a5
commit ab71643c99
23 changed files with 82 additions and 52 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/group.cc
+6 -2
View File
@@ -202,7 +202,7 @@ ncclResult_t ncclGroupEnd() {
if (args->funcType == ASYNC_FUNC_COLL && args->coll.comm->connect) {
int err = pthread_join(ncclGroupThreads[i], NULL);
if (err != 0) {
WARN("Error waiting for pthread_join : %s\n", strerror(errno));
WARN("Error waiting for pthread_join : %s", strerror(errno));
return ncclSystemError;
}
NCCLCHECKGOTO(args->ret, ret, end);
@@ -234,7 +234,7 @@ ncclResult_t ncclGroupEnd() {
// schedule delta 0, +1, -1, +2, -2, ...
// also make sure we don't do 0 twice, nor +n/2 and -n/2 if n is even.
for (int d=0; d<=nRanks/4; d++) {
int deltas[4] = { d, (nRanks-d)%nRanks, nRanks/2-d, nRanks-(nRanks/2-d) };
int deltas[4] = { d, (nRanks-d)%nRanks, nRanks/2-d, (nRanks-(nRanks/2-d))%nRanks };
int index = 0;
int delta = deltas[index];
sched_delta:
@@ -259,6 +259,10 @@ sched_delta:
ssize_t sendbytes = totSendBytes-sendOffset;
if (recvbytes > recvChunkSize) { recvbytes = recvChunkSize; } else { recvRemaining = 0; }
if (sendbytes > sendChunkSize) { sendbytes = sendChunkSize; } else { sendRemaining = 0; }
// 0-bytes send/recv are considered as syncs. Make sure we only add syncs when requested
// (total size == 0), otherwise set size to -1 so that the kernel skips the operation.
if (sendbytes == 0 && totSendBytes != 0) sendbytes = -1;
if (recvbytes == 0 && totRecvBytes != 0) recvbytes = -1;
if (sendbytes >= 0 || recvbytes >= 0) {
NCCLCHECKGOTO(scheduleSendRecv(comm, delta, channelId,
recvbytes, recv ? ((char*)(recv->buff)) + recvOffset : NULL,