Welcome to mirror list, hosted at ThFree Co, Russian Federation.

github.com/marian-nmt/nccl.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorSylvain Jeaugey <sjeaugey@nvidia.com>2018-11-20 04:43:50 +0300
committerSylvain Jeaugey <sjeaugey@nvidia.com>2018-11-27 03:24:31 +0300
commit98adf2fe11bb13f93e93c805ffcb28fd9f0578a1 (patch)
treec45afebf0c856bfecb27165d740902c1ce5eb0de /src/transport/net.cu
parent0d3a20f96d4887bee86a0fd7bf79feb14e5a01f5 (diff)
Make network isend/irecv non blocking
Diffstat (limited to 'src/transport/net.cu')
-rw-r--r--src/transport/net.cu20
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;