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 diff --git a/projects/rccl/src/misc/utils.cc b/projects/rccl/src/misc/utils.cc index 74190de6b7..eed57140df 100644 --- a/projects/rccl/src/misc/utils.cc +++ b/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= '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; } diff --git a/projects/rccl/src/transport.cc b/projects/rccl/src/transport.cc index e66a6f25ef..62940498cb 100644 --- a/projects/rccl/src/transport.cc +++ b/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) { diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc index be24232ca9..be459840ca 100644 --- a/projects/rccl/src/transport/net.cc +++ b/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; }