diff options
author | Sylvain Jeaugey <sjeaugey@nvidia.com> | 2018-11-20 04:43:50 +0300 |
---|---|---|
committer | Sylvain Jeaugey <sjeaugey@nvidia.com> | 2018-11-27 03:24:31 +0300 |
commit | 98adf2fe11bb13f93e93c805ffcb28fd9f0578a1 (patch) | |
tree | c45afebf0c856bfecb27165d740902c1ce5eb0de /src/transport/net.cu | |
parent | 0d3a20f96d4887bee86a0fd7bf79feb14e5a01f5 (diff) |
Make network isend/irecv non blocking
Diffstat (limited to 'src/transport/net.cu')
-rw-r--r-- | src/transport/net.cu | 20 |
1 files changed, 13 insertions, 7 deletions
diff --git a/src/transport/net.cu b/src/transport/net.cu index 8a7e3b8..165187c 100644 --- a/src/transport/net.cu +++ b/src/transport/net.cu @@ -416,17 +416,21 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) { while (f1[0] != flag || f2[0] != flag); } NCCLCHECK(ncclNetIsend(resources->netSendComm, lines, size, ptrType, requests+slot)); - sizesFifo[slot] = size; - tail++; - idle = 0; + if (requests[slot] != NULL) { + sizesFifo[slot] = size; + tail++; + idle = 0; + } } } } else while (tail < *prevTail) { // Send through network int slot = tail%args->substeps; NCCLCHECK(ncclNetIsend(resources->netSendComm, localBuff+slot*sliceSize, sizesFifo[slot], ptrType, requests+slot)); - tail++; - idle = 0; + if (requests[slot] != NULL) { + tail++; + idle = 0; + } } if (head < tail) { int done; @@ -502,8 +506,10 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) { if ((tail < head + args->substeps) && (tail < *nextHead + args->substeps) && (tail < end)) { int slot = tail%args->substeps; NCCLCHECK(ncclNetIrecv(resources->netRecvComm, localBuff+slot*sliceSize, sliceSize, ptrType, requests+slot)); - tail++; - idle = 0; + if (requests[slot] != NULL) { + tail++; + idle = 0; + } } if (tail > head) { int done; |