diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h
index 81a4d4cb7f..8c63a6a4f1 100644
--- a/src/collectives/device/primitives.h
+++ b/src/collectives/device/primitives.h
@@ -193,10 +193,11 @@ class ncclPrimitives {
       if (SEND) __threadfence_system();
       if (tid == 0) FOR_SEND(postSend);
       if (tid == 0) FOR_RECV(postRecv);
+
+      for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
+      for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;
+      offset += sliceSize;
     }
-    for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
-    for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;
-    offset += sliceSize;
   }
 
   __device__ void loadRecvConn(struct ncclConnInfo* conn, int i, T* directBuff) {