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, */