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

[gfx950] Make bypassing __threadfence the default for multinode. (#1947)

Browse Source 
* Gate based on ROCM version, safe for ROCm 7.0.2 and beyond.
* Updates naming to gfx9CheapFenceOff since we use this for gfx942 and gfx950.  Thanks Nilesh.
* Add info logging statement to NCCL_INIT to print whether enabled when INFO logging is enabled.
This commit is contained in:
alex-breslow-amd
2025-10-15 09:15:36 -07:00
committed by GitHub
parent 0f99fd84a3
commit c70f5b4621
5 changed files with 12 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/enqueue.cc
+2 -2
View File
@@ -348,7 +348,7 @@ static bool testBudget(
// Returns whether this should be disabled at the device level. Should be called after devWork fields have been set for what
// it depends on.
bool gfx942CheapFenceOff(const ncclDevWorkColl& devWork, bool disabledByPrecheck){
bool gfx9CheapFenceOff(const ncclDevWorkColl& devWork, bool disabledByPrecheck){
bool fenceOk = devWork.regUsed == 0 && devWork.netRegUsed == 0 && !disabledByPrecheck;
return !fenceOk;
}
@@ -388,7 +388,7 @@ ncclResult_t ncclTasksRegAndEnqueue(struct ncclComm* comm) {
devWork.isOneRPN = comm->isOneRPN;
devWork.netRegUsed = devWork.regUsed = 0;
devWork.gfx942CheapFenceOff = gfx942CheapFenceOff(devWork, comm->gfx942CheapFenceOff);
devWork.gfx9CheapFenceOff = gfx9CheapFenceOff(devWork, comm->gfx9CheapFenceOff);
devWork.profilerEnabled = ncclProfilerPluginLoaded() && (task->eActivationMask & ncclProfileKernelCh);
if (task->regBufType & NCCL_NET_REG_BUFFER)
devWork.netRegUsed = 1;