Allow zero byte sendrecv in alltoallv (#1349)
* Allow zero byte sendrecv in alltoallv * Fix previous merge error
This commit is contained in:
@@ -142,14 +142,14 @@ ncclResult_t ncclAllToAllv_impl(const void *sendbuff, const size_t sendcounts[],
|
||||
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
||||
NCCLCHECK(ncclGroupStart());
|
||||
for (int r=0; r<nRanks; r++) {
|
||||
if (sendcounts[r]) NCCLCHECK(ncclSend(
|
||||
NCCLCHECK(ncclSend(
|
||||
((char*)sendbuff) + sdispls[r]*ncclTypeSize(datatype),
|
||||
sendcounts[r],
|
||||
datatype,
|
||||
r,
|
||||
comm,
|
||||
stream));
|
||||
if (recvcounts[r]) NCCLCHECK(ncclRecv(
|
||||
NCCLCHECK(ncclRecv(
|
||||
((char*)recvbuff) + rdispls[r]*ncclTypeSize(datatype),
|
||||
recvcounts[r],
|
||||
datatype,
|
||||
|
||||
+2
-2
@@ -609,7 +609,7 @@ static ncclResult_t addP2pToPlan(
|
||||
|
||||
struct ncclWorkElemP2p elem = {0};
|
||||
elem.proto = info.protocol;
|
||||
elem.peer = addr == nullptr ? -1 : peer;
|
||||
elem.peer = peer;
|
||||
elem.nWarps = NCCL_MAX_NTHREADS/comm->WarpSize;
|
||||
elem.reg = proxyOp.reg;
|
||||
elem.p2pType = isSendNotRecv ? ncclWorkP2pTypeSend : ncclWorkP2pTypeRecv;
|
||||
@@ -628,7 +628,7 @@ static ncclResult_t addP2pToPlan(
|
||||
// Calculate the opCount after appendWorkElemP2p since it will always return
|
||||
// with channel->nWork equal to one plus the work index this p2p settled in.
|
||||
proxyOp.opCount = uint64_t(plan->channels[channelId].nWork)<<1 | 1;
|
||||
if (addr != nullptr) NCCLCHECK(addProxyOpIfNeeded(comm, plan, &proxyOp));
|
||||
NCCLCHECK(addProxyOpIfNeeded(comm, plan, &proxyOp));
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
|
||||
Viittaa uudesa ongelmassa
Block a user