NCCL 2.4.6-1

Added detection of IBM/Power NVLink bridge device.
    Add NUMA support to PCI distance calculations.
    Added NCCL_IGNORE_CPU_AFFINITY env var.
    Fix memory leaks; GithubIssue#180
    Compiler warning fix; GithubIssue#178
    Replace non-standard variable length arrays. GithubIssue#171
    Fix Tree+Shared Memory crash. GithubPR#185
    Fix LL cleanup hang during long running DL jobs.
    Fix NCCL_RINGS environment variable handling.
    Added extra checks to catch repeat calls to ncclCommDestroy() GithubIssue#191
    Improve bootstrap socket connection reliability at scale.
    Fix hostname hashing issue. GithubIssue#187
    Code cleanup to rename all non device files from *.cu to *.cc


[ROCm/rccl commit: f40ce73e89]

projects/rccl/src/misc/group.cc

@@ -0,0 +1,198 @@
+/*************************************************************************
+ * Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "group.h"
+#include "debug.h"
+#include "enqueue.h"
+
+#define MAX_ASYNC_OPS 128
+thread_local pthread_t ncclGroupThreads[MAX_ASYNC_OPS];
+thread_local int ncclGroupIndex = 0;
+thread_local int ncclGroupMode = 0;
+thread_local ncclResult_t ncclGroupError = ncclSuccess;
+
+bool ncclAsyncMode() {
+  return ncclGroupMode > 0;
+}
+
+ncclResult_t ncclAsyncErrCheck(ncclResult_t ret) {
+  if (ncclGroupError == ncclSuccess || ret != ncclSuccess) ncclGroupError = ret;
+  return ret;
+}
+
+struct ncclInitArgs {
+  ncclInitFunc_t func;
+  int cudaDev;
+  ncclComm_t* newcomm;
+  int ndev;
+  ncclUniqueId commId;
+  int myrank;
+};
+struct ncclCollArgs {
+  ncclComm_t comm;
+};
+
+enum ncclAsyncFuncType {
+  ASYNC_FUNC_INVALID = 0,
+  ASYNC_FUNC_INIT = 1,
+  ASYNC_FUNC_COLL = 2,
+};
+struct ncclAsyncArgs {
+  ncclResult_t ret;
+  enum ncclAsyncFuncType funcType;
+  union {
+    ncclCollArgs coll;
+    ncclInitArgs init;
+  };
+};
+
+thread_local struct ncclAsyncArgs ncclGroupArgs[MAX_ASYNC_OPS];
+
+ncclResult_t ncclSetDevice(int cudaDev) {
+  CUDACHECK(cudaSetDevice(cudaDev));
+  return ncclSuccess;
+}
+
+#define CHECK(a) do { \
+  if ((args->ret = (a)) != ncclSuccess) { \
+    INFO(NCCL_INIT,"%s:%d -> %d [Async thread]", __FILE__, __LINE__, args->ret); \
+    return args; \
+  } \
+} while(0)
+
+void* ncclAsyncThreadMain(void* args_) {
+  struct ncclAsyncArgs* args = (struct ncclAsyncArgs*)args_;
+  CHECK(ncclSetDevice(args->init.cudaDev));
+  CHECK(args->init.func(args->init.newcomm, args->init.ndev, args->init.commId, args->init.myrank));
+  return args;
+}
+
+ncclResult_t ncclAsyncInit(ncclInitFunc_t func, int cudaDev, ncclComm_t* newcomm, int ndev, ncclUniqueId commId, int myrank) {
+  if (ncclGroupIndex >= MAX_ASYNC_OPS) {
+    WARN("Too many async operations in progress, max is %d", MAX_ASYNC_OPS);
+    return ncclAsyncErrCheck(ncclInternalError);
+  }
+  int index = ncclGroupIndex++;
+  struct ncclAsyncArgs* args = ncclGroupArgs+index;
+  args->funcType = ASYNC_FUNC_INIT;
+  args->init.func = func;
+  args->init.cudaDev = cudaDev;
+  args->init.newcomm = newcomm;
+  args->init.ndev = ndev;
+  memcpy(&args->init.commId, &commId, sizeof(commId));
+  args->init.myrank = myrank;
+  // We need to use threads for Init
+  pthread_create(ncclGroupThreads+index, NULL, ncclAsyncThreadMain, args);
+  return ncclSuccess;
+}
+
+ncclResult_t ncclAsyncColl(ncclComm_t comm) {
+  struct ncclAsyncArgs* args = ncclGroupArgs;
+  for (int i=0; i<ncclGroupIndex; i++) {
+    if (args->coll.comm == comm) return ncclSuccess;
+    args++;
+  }
+  if (ncclGroupIndex >= MAX_ASYNC_OPS) {
+    WARN("Too many async operations in progress, max is %d", MAX_ASYNC_OPS);
+    return ncclAsyncErrCheck(ncclInternalError);
+  }
+  ncclGroupIndex++;
+  args->funcType = ASYNC_FUNC_COLL;
+  args->coll.comm = comm;
+  return ncclSuccess;
+}
+
+NCCL_API(ncclResult_t, ncclGroupStart);
+ncclResult_t ncclGroupStart() {
+  ncclGroupMode++;
+  return ncclSuccess;
+}
+
+NCCL_API(ncclResult_t, ncclGroupEnd);
+ncclResult_t ncclGroupEnd() {
+  ncclGroupMode--;
+  if (ncclGroupMode > 0) return ncclSuccess;
+  int savedDev;
+  CUDACHECK(cudaGetDevice(&savedDev));
+  int done = ncclGroupIndex;
+  int doneArray[MAX_ASYNC_OPS];
+  for (int i=0; i<ncclGroupIndex; i++) doneArray[i] = 0;
+
+  ncclResult_t ret = ncclGroupError;
+  if (ret != ncclSuccess) goto group_cleanup;
+
+  /* Collectives are done in three steps :
+   * 1. Barrier Check In. Only the last call may call cudaLaunchKernel[cooperative]
+   * 2. Barrier Wait. No CUDA call is permitted
+   * 3. Enqueue Events. CUDA event wait/enqueue.
+   * This is needed because step 2 cannot call any CUDA primitive, otherwise if
+   * cudaFree happens between 1 and 3, it could block that CUDA call and
+   * prevent some ranks from launching their network threads, which would
+   * prevent the NCCL call from completing, blocking the cudaFree call.
+   */
+  for (int i=0; i<ncclGroupIndex; i++) {
+    struct ncclAsyncArgs* args = ncclGroupArgs+i;
+    if (args->funcType == ASYNC_FUNC_COLL) {
+      if (args->coll.comm->userStream == NULL)
+        CUDACHECKGOTO(cudaSetDevice(args->coll.comm->cudaDev), ret, end);
+      NCCLCHECKGOTO(ncclBarrierEnqueue(args->coll.comm), ret, end);
+    }
+  }
+  for (int i=0; i<ncclGroupIndex; i++) {
+    struct ncclAsyncArgs* args = ncclGroupArgs+i;
+    if (args->funcType == ASYNC_FUNC_COLL) {
+      CUDACHECKGOTO(cudaSetDevice(args->coll.comm->cudaDev), ret, end);
+      NCCLCHECKGOTO(ncclBarrierEnqueueWait(args->coll.comm), ret, end);
+    }
+  }
+  for (int i=0; i<ncclGroupIndex; i++) {
+    struct ncclAsyncArgs* args = ncclGroupArgs+i;
+    if (args->funcType == ASYNC_FUNC_COLL) {
+      if (args->coll.comm->userStream == NULL)
+        CUDACHECKGOTO(cudaSetDevice(args->coll.comm->cudaDev), ret, end);
+      NCCLCHECKGOTO(ncclEnqueueEvents(args->coll.comm), ret, end);
+      doneArray[i] = 1;
+      done--;
+    }
+  }
+
+  /* For init, since we use threads, we just wait for threads to complete */
+  while (done) {
+    for (int i=0; i<ncclGroupIndex; i++) {
+      struct ncclAsyncArgs* args = ncclGroupArgs+i;
+      if (args->funcType == ASYNC_FUNC_INIT && doneArray[i] == 0) {
+        int err = pthread_tryjoin_np(ncclGroupThreads[i], NULL);
+        if (err == EBUSY) continue;
+        if (err != 0) { ret = ncclSystemError; goto end; }
+        if (args->ret != ncclSuccess) { ret = args->ret; goto end; }
+        doneArray[i] = 1;
+        done--;
+      }
+    }
+  }
+  goto end;
+group_cleanup:
+  // At least one call in the group failed. Since we want to make that group
+  // an atomic operation, we need to cancel all operations.
+  for (int i=0; i<ncclGroupIndex; i++) {
+    struct ncclComm* comm = ncclGroupArgs[i].coll.comm;
+    for (int c=0; c<comm->nChannels; c++) {
+      struct ncclChannel* channel = comm->channels+c;
+      for (int i=0; i<channel->collCount; i++) {
+        channel->collectives[(channel->collStart + i)%NCCL_MAX_OPS].active = 0;
+      }
+      channel->collFifoTail = channel->collStart;
+      channel->collCount = 0;
+    }
+    comm->myParams->gridDim.x = comm->myParams->blockDim.x = 0;
+    comm->userStreamSet = false;
+  }
+end:
+  ncclGroupError = ncclSuccess;
+  ncclGroupIndex = 0;
+  CUDACHECK(cudaSetDevice(savedDev)); // do other clean-ups first before calling cudaSetDevice, because this call can fail too
+  return ret;
+}