From a52ca36a9eeb1d1d7738f4d4cd0ed43f9e563323 Mon Sep 17 00:00:00 2001
From: Chris Jones <cjfj@deepmind.com>
Date: Thu, 18 Nov 2021 14:21:05 +0000
Subject: [PATCH 1/3] Perform `busIdToInt64` on the stack.

I noticed when I enabled `NCCL_DEBUG_SUBSYS=ALLOC` that this function is
called thousands of times, making the log output unintelligible.
Fortunately, this function can be implemented without heap allocations.


[ROCm/rccl commit: 8cf7325d69184d01bd37fb50d03f8a05e647c5b6]
---
 projects/rccl/src/misc/utils.cc | 7 ++-----
 1 file changed, 2 insertions(+), 5 deletions(-)

diff --git a/projects/rccl/src/misc/utils.cc b/projects/rccl/src/misc/utils.cc
index 79e61704dd..f3e3ca264e 100644
--- a/projects/rccl/src/misc/utils.cc
+++ b/projects/rccl/src/misc/utils.cc
@@ -25,11 +25,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') ||
@@ -40,7 +38,6 @@ ncclResult_t busIdToInt64(const char* busId, int64_t* id) {
   }
   hexStr[hexOffset] = '\0';
   *id = strtol(hexStr, NULL, 16);
-  free(hexStr);
   return ncclSuccess;
 }
 

From 07d9648aec2b1e5a36c5597089ee0d0e99484ff9 Mon Sep 17 00:00:00 2001
From: Ke Wen <kw2501@fb.com>
Date: Tue, 23 Nov 2021 14:43:14 -0800
Subject: [PATCH 2/3] Add env NCCL_NET_DISABLE_INTRA

Disable NET transport for intra-node communication by setting the env to 1
It provides an option to error out instead of falling back to NET when superior intra-node transports (P2P and SHM) are unavailable


[ROCm/rccl commit: c88c9f873fa680b547c4ae177c8f915d5c2dd157]
---
 projects/rccl/src/transport.cc     |  4 ++--
 projects/rccl/src/transport/net.cc | 10 ++++++++++
 2 files changed, 12 insertions(+), 2 deletions(-)

diff --git a/projects/rccl/src/transport.cc b/projects/rccl/src/transport.cc
index d7eadcdd87..2cb5538a69 100644
--- a/projects/rccl/src/transport.cc
+++ b/projects/rccl/src/transport.cc
@@ -36,8 +36,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) {
diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc
index 2b548cee39..5abc32d15b 100644
--- a/projects/rccl/src/transport/net.cc
+++ b/projects/rccl/src/transport/net.cc
@@ -56,8 +56,18 @@ struct netRecvResources {
   uint64_t llLastCleaning;
 };
 
+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;
 }

From 04f51513a64141495f1197c73c971b6e9ab4826d Mon Sep 17 00:00:00 2001
From: Chang Lan <changlan@google.com>
Date: Wed, 8 Dec 2021 11:43:36 -0800
Subject: [PATCH 3/3] Build fastsocket plugin from ext-net

[ROCm/rccl commit: c5790b36722d5b41ee2a9b2bad69e364180ffd22]
---
 .../rccl/ext-net/google-fastsocket/Makefile   | 22 +++++++++++++++++++
 1 file changed, 22 insertions(+)
 create mode 100644 projects/rccl/ext-net/google-fastsocket/Makefile

diff --git a/projects/rccl/ext-net/google-fastsocket/Makefile b/projects/rccl/ext-net/google-fastsocket/Makefile
new file mode 100644
index 0000000000..e40e3053ad
--- /dev/null
+++ b/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