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

gfx950 channel tuning for ReduceScatter and AllGather (#1940)

Browse Source 
* add channel thresholds to override channel-count adjustments
This commit is contained in:
isaki001
2025-10-14 09:50:44 -05:00
committed by GitHub
parent e738c03e39
commit 0f99fd84a3
7 changed files with 88 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/enqueue.cc
+9 -1
View File
@@ -2015,6 +2015,8 @@ static ncclResult_t updateCollCostTable(
return ncclSuccess;
}
extern int64_t ncclParamMinNchannels();
static ncclResult_t topoGetAlgoInfo(
struct ncclComm* comm, struct ncclTaskColl* info, size_t nBytes,
float** collCostTable, ncclSimInfo_t* simInfo
@@ -2079,11 +2081,17 @@ static ncclResult_t topoGetAlgoInfo(
nc = comm->nvlsChannels;
} else {
rcclUpdateThreadThreshold(comm, nBytes, info, threadThreshold);
INFO(NCCL_INIT, "pre-adjustment threadThreshold:%i nBytes:%lu nc:%i", threadThreshold, nBytes, nc);
int minNChannels = ncclParamMinNchannels();
// Ring/Tree channel tuning
while (nBytes < nc * nt * threadThreshold) {
INFO(NCCL_INIT, "minNChannels:%i", minNChannels);
while (nBytes < nc * nt * threadThreshold && nc > minNChannels) {
if (nc >= 2) nc--;
else break;
}
INFO(NCCL_INIT, "post-adjustment based on threadThreshold:%i nBytes:%lu nc:%i", threadThreshold, nBytes, nc);
rcclOverrideChannels(comm, info->func, nBytes, nc);
}
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
#else