AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fix MSCCL multi-node (#1032)

Browse Source 
1) Move needsProxy initialization before mscclSetupConnections since the latter
will revise it later.
2) Remove mscclAvailable check in net.cc since it's no more required and caused
non-shared buffer allocated for MSCCL which is not expected.
This commit is contained in:
Ziyue Yang
2024-01-06 09:03:43 +08:00
committed by GitHub
parent cfc04a8aef
commit 70bbeb4773
2 changed files with 5 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/misc/msccl/msccl_lifecycle.cc
+3 -3
View File
@@ -211,12 +211,13 @@ ncclResult_t mscclInit(ncclComm_t comm) {
mscclStatus& status = mscclGetStatus();
// Free algorithm handles are initialized globally once and before algorithm pre-processing
// freeAlgoHandles and needsProxy are initialized globally once and before algorithm pre-processing and connection
if (!mscclInitialized.load(std::memory_order_acquire)) {
status.freeAlgoHandles.resize(MSCCL_MAX_NUM_ALGOS);
for (int i = 0; i < MSCCL_MAX_NUM_ALGOS; i++) {
status.freeAlgoHandles[i] = MSCCL_MAX_NUM_ALGOS - i - 1;
}
status.needsProxy = false;
}
// Pre-process all algorithms for internal scheduler and for different comms.
@@ -231,7 +232,7 @@ ncclResult_t mscclInit(ncclComm_t comm) {
NCCLCHECK(mscclLoadAlgo(m.filePath.c_str(), &(status.rankToAlgoHandles[i][comm->rank]), comm->rank));
}
// Connect algorithms
mscclAlgoHandle_t mscclAlgoHandle = status.rankToAlgoHandles[i][comm->rank];
mscclAlgoHandle_t mscclAlgoHandle = status.rankToAlgoHandles[i][comm->rank];
if (status.connectedAlgos[comm].find(mscclAlgoHandle) == status.connectedAlgos[comm].end()) {
NCCLCHECK(mscclSetupConnections(status.hostAlgos[mscclAlgoHandle], comm));
status.connectedAlgos[comm].insert(mscclAlgoHandle);
@@ -247,7 +248,6 @@ ncclResult_t mscclInit(ncclComm_t comm) {
status.workIndex = 1;
NCCLCHECK(ncclCudaCalloc(&status.syncFlags, MSCCL_MAX_NUM_THREAD_BLOCKS));
status.lastStream = nullptr;
status.needsProxy = false;
NCCLCHECK(mscclInitWorkFifoStatus(&(status.defaultWorkFifoStatus)));
mscclInitialized.store(true, std::memory_order_release);
src/transport/net.cc
+2 -2
View File
@@ -179,7 +179,7 @@ static ncclResult_t sendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph
struct setupReq req = { 0 };
int localRank, tpProxyRank;
send->conn.shared = req.shared = (graph || mscclAvailable() && mscclIsCaller()) ? 0 : ncclParamNetSharedBuffers() != -2 ? ncclParamNetSharedBuffers() : 1;
send->conn.shared = req.shared = (graph || mscclIsCaller()) ? 0 : ncclParamNetSharedBuffers() != -2 ? ncclParamNetSharedBuffers() : 1;
req.channelId = channelId;
req.connIndex = connIndex;
req.curr_hdp_reg = 0;
@@ -224,7 +224,7 @@ static ncclResult_t recvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph
struct setupReq req = { 0 };
int localRank;
recv->conn.shared = req.shared = (graph || mscclAvailable() && mscclIsCaller()) ? 0 : ncclParamNetSharedBuffers() != -2 ? ncclParamNetSharedBuffers() : 1;
recv->conn.shared = req.shared = (graph || mscclIsCaller()) ? 0 : ncclParamNetSharedBuffers() != -2 ? ncclParamNetSharedBuffers() : 1;
req.channelId = channelId;
req.connIndex = connIndex;
req.netDev = -1;