Use One Slice per Basic Primitive for AllReduce, ReduceScatter, AllGather (#1681) for Single Node on Some GFX9 Systems

Using a single slice rather than the typical two provides about 5% speedup (sometimes more or less) on some GFX9 systems for single node.

[ROCm/rccl commit: 2f6b20c00a]
Tento commit je obsažen v:
alex-breslow-amd authored and GitHub committed 2025-05-29 16:17:35 -07:00
1 rodič 19ed482121
revize 4277b5aa88
9 files changed
+66 -10

No files matched your search

+1
Zobrazit soubor
@@ -332,6 +332,7 @@ ncclResult_t ncclTasksRegAndEnqueue(struct ncclComm* comm) {
devWork.redOpArg = task->opDev.scalarArg;
devWork.redOpArgIsPtr = task->opDev.scalarArgIsPtr;
devWork.oneNode = (comm->nNodes == 1);
devWork.rcclUseOneSlice = comm->rcclUseOneSlice;
devWork.isOneRPN = comm->isOneRPN;
devWork.netRegUsed = devWork.regUsed = 0;
if (task->regBufType & NCCL_NET_REG_BUFFER)