Merge remote-tracking branch 'nccl/master' into develop

[ROCm/rccl commit: 973d0111db]

projects/rccl/ext-net/google-fastsocket/Makefile

@@ -0,0 +1,22 @@
+CUDA_HOME?=/usr/local/cuda
+INC:=-I$(CUDA_HOME)/include
+PLUGIN_SO:=../../build/libnccl-net.so
+
+default: $(PLUGIN_SO)
+
+$(PLUGIN_SO): nccl-fastsocket/net_fastsocket.cc nccl-fastsocket/compat.cc
+	$(CC) $(INC) -fPIC -shared -o $@ -Wl,-soname,$(PLUGIN_SO) $^
+
+nccl-fastsocket/%.cc:
+	git clone https://github.com/google/nccl-fastsocket.git
+
+install: $(BUILDDIR)/lib/$(PLUGIN_SO)
+
+$(BUILDDIR)/lib/$(PLUGIN_SO): $(PLUGIN_SO)
+	@printf "Grabbing %-35s > %s\n" $< $@
+	mkdir -p $(BUILDDIR)/lib
+	install -m 644 $< $@
+
+clean:
+	rm -f $(PLUGIN_SO)
+	rm -Rf nccl-fastsocket

projects/rccl/src/misc/utils.cc

@@ -27,11 +27,9 @@ ncclResult_t int64ToBusId(int64_t id, char* busId) {
 }
 
 ncclResult_t busIdToInt64(const char* busId, int64_t* id) {
-  const int size = strlen(busId);
-  char* hexStr;
-  NCCLCHECK(ncclCalloc(&hexStr, size));
+  char hexStr[17];  // Longest possible int64 hex string + null terminator.
   int hexOffset = 0;
-  for (int i=0; i<size; i++) {
+  for (int i = 0; hexOffset < sizeof(hexStr) - 1; i++) {
     char c = busId[i];
     if (c == '.' || c == ':') continue;
     if ((c >= '0' && c <= '9') ||
@@ -42,7 +40,6 @@ ncclResult_t busIdToInt64(const char* busId, int64_t* id) {
   }
   hexStr[hexOffset] = '\0';
   *id = strtol(hexStr, NULL, 16);
-  free(hexStr);
   return ncclSuccess;
 }
 

projects/rccl/src/transport.cc

@@ -49,8 +49,8 @@ static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph*
       return ncclSuccess;
     }
   }
-  WARN("No transport found !");
-  return ncclInternalError;
+  WARN("No transport found for rank %d[%lx] -> rank %d[%lx]", myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId);
+  return ncclSystemError;
 }
 
 ncclResult_t ncclTransportP2pConnect(struct ncclComm* comm, struct ncclChannel* channel, int nrecv, int* peerRecv, int nsend, int* peerSend, int connIndex) {

projects/rccl/src/transport/net.cc

@@ -61,8 +61,18 @@ struct netRecvResources {
   uint32_t* curr_hdp_reg;  // Curr GPU in ring (for rdma transport use only)
 };
 
+NCCL_PARAM(NetDisableIntra, "NET_DISABLE_INTRA", -2);
+
 /* Determine if two peers can communicate with NET */
 ncclResult_t netCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
+  // Same host?
+  if (info1->hostHash == info2->hostHash) {
+    // User disabled NET for intra-node?
+    if (ncclParamNetDisableIntra() == 1) {
+      *ret = 0;
+      return ncclSuccess;
+    }
+  }
   *ret = 1;
   return ncclSuccess;
 }