yezhengmao/nccl
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fix Collnet when GDR is disabled

Browse Source
This commit is contained in:
Sylvain Jeaugey 2021-09-22 05:10:09 -07:00
parent e11238b302
commit 4ec992fab7
1 changed files with 1 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/transport/coll_net.cc
View File

@ -507,7 +507,7 @@ ncclResult_t collNetRecvProxy(struct ncclProxyArgs* args) {
int sharedBuffSlot = sub->transmitted%NCCL_STEPS; int sharedBuffSlot = sub->transmitted%NCCL_STEPS;
int startChannel = group*COLLNET_GROUP_NSUBS; int startChannel = group*COLLNET_GROUP_NSUBS;
char* groupRecvAddress; char* groupRecvAddress;
NCCLCHECK(ncclProxySharedBuffersGetCollNet(sub->connector->comm, 1, 1, sharedBuffSlot, startChannel, &groupRecvAddress)); NCCLCHECK(ncclProxySharedBuffersGetCollNet(sub->connector->comm, p == NCCL_PROTO_SIMPLE ? resources->useGdr : 0, 1, sharedBuffSlot, startChannel, &groupRecvAddress));
char* ptr = groupRecvAddress + (s%COLLNET_GROUP_NSUBS)*args->sharedSize[sharedBuffSlot]; char* ptr = groupRecvAddress + (s%COLLNET_GROUP_NSUBS)*args->sharedSize[sharedBuffSlot];
if (p == NCCL_PROTO_SIMPLE) { if (p == NCCL_PROTO_SIMPLE) {
volatile void** ptrsFifo = (volatile void**)resources->recvMem->ptrsFifo; volatile void** ptrsFifo = (volatile void**)resources->recvMem->ptrsFifo;