From 4180e6409e79c381cfe18a6b98d80e79dbb8b30b Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Mon, 31 Aug 2020 17:33:22 -0700 Subject: [PATCH] Fix incorrect threads split in sendrecv (#261) --- src/collectives/device/sendrecv.h | 10 +--------- 1 file changed, 1 insertion(+), 9 deletions(-) diff --git a/src/collectives/device/sendrecv.h b/src/collectives/device/sendrecv.h index 8dab76d6ad..ff1e4e75f0 100644 --- a/src/collectives/device/sendrecv.h +++ b/src/collectives/device/sendrecv.h @@ -40,15 +40,7 @@ __device__ void ncclSendRecvKernel(struct CollectiveArgs* args) { const int stepSize = comm->buffSizes[NCCL_PROTO_SIMPLE]/(sizeof(T)*NCCL_STEPS)/SENDRECV_SLICEFACTOR; - int nthreadsSplit; - const int64_t sendCount = static_cast(args->p2p.sendCount); - const int64_t recvCount = static_cast(args->p2p.recvCount); - if (sendCount >= 0 && recvCount >= 0) - nthreadsSplit = nthreads/2; - else { - if (sendCount >= 0) nthreadsSplit = nthreads; - else nthreadsSplit = 0; - } + int nthreadsSplit = nthreads/2; // We set NRECV or NSEND to 2 to use different barriers in primitives for the send threads and // receive threads, but then we define all peers to -1 since sender threads don't receive and // receive threads don't send.