From b27490d38d92c41b277c2134878850f5228549a3 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Wed, 19 May 2021 08:25:26 -0700 Subject: [PATCH] Use fixed segment size for sendrecv (#369) [ROCm/rccl commit: 8cde34be51b16e56f052095cd2ab7d6c6107d11f] --- projects/rccl/src/enqueue.cc | 8 +++----- projects/rccl/src/include/devcomm.h | 2 +- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/projects/rccl/src/enqueue.cc b/projects/rccl/src/enqueue.cc index b7315678ff..e214212f6f 100644 --- a/projects/rccl/src/enqueue.cc +++ b/projects/rccl/src/enqueue.cc @@ -714,8 +714,8 @@ static ncclResult_t ncclSaveP2p(struct ncclInfo* info) { return ncclSuccess; } -static int getSegment(int delta, struct ncclWork* work, int e) { - for (int s=0; selems[s].p2p.delta != delta; s++) { +static int getSegment(int delta, struct ncclWork* work) { + for (int s=0; selems[s].p2p.delta != delta; s++) { if (work->elems[s].p2p.nThreads == 0) return s; } return -1; @@ -759,11 +759,9 @@ ncclResult_t ncclEnqueueP2pKernel(struct ncclComm* comm, struct ncclQueueElem* e int opIndex = (channel->workFifoTail-1+NCCL_MAX_OPS)%NCCL_MAX_OPS; struct ncclWork* w = channel->workFifo+opIndex; int segment = -1; - const int e = ((comm->topo->nodes[GPU].count == comm->topo->nRanks) && (comm->topo->type & RCCL_TOPO_4P2H_ROME)) - ? 1 : NCCL_MAX_WORK_ELEMENTS; if (channel->workCount && w->elems[0].funcIndex == FUNC_INDEX_P2P && w->elems[NCCL_MAX_WORK_ELEMENTS-1].p2p.nThreads == 0) { // Try to pack more segments into a single operation - segment = getSegment(workElem->p2p.delta, w, e); + segment = getSegment(workElem->p2p.delta, w); } if (segment == -1) { NCCLCHECK(getNextOp(channel, &w, NULL)); diff --git a/projects/rccl/src/include/devcomm.h b/projects/rccl/src/include/devcomm.h index 0221e1e99b..4c7c9980bc 100644 --- a/projects/rccl/src/include/devcomm.h +++ b/projects/rccl/src/include/devcomm.h @@ -175,7 +175,7 @@ struct ncclDevComm; #pragma pack(push) /* push current alignment to stack */ #pragma pack(4) /* set alignment to 4 bytes boundary */ -#define NCCL_MAX_WORK_ELEMENTS 2 +#define NCCL_MAX_WORK_ELEMENTS 1 #define NCCL_MAX_GROUPS (NCCL_MAX_WORK_ELEMENTS*2) /* ncclWork is to be a power of two, currently 8x64 bytes, */