diff --git a/src/clique/AllReduceCliqueKernel.h b/src/clique/AllReduceCliqueKernel.h
index 5fc4752efb..d941a6c045 100644
--- a/src/clique/AllReduceCliqueKernel.h
+++ b/src/clique/AllReduceCliqueKernel.h
@@ -66,7 +66,7 @@ __device__ void AllReduceCliqueSplitKernel(struct ncclWorkElem* args)
     // Perform the reduction
     #define ALL_REDUCE_CLIQUE_UNROLL 1
     ReduceOrCopyMulti<ALL_REDUCE_CLIQUE_UNROLL, FUNC, T, NUM_RANKS, NUM_RANKS, NUM_RANKS, NUM_RANKS>(
-      threadIdx.x, blockDim.x, redOp, true, true, NUM_RANKS, srcs, NUM_RANKS, dsts, blockN);
+      threadIdx.x, blockDim.x, redOp, NUM_RANKS, true, NUM_RANKS, srcs, NUM_RANKS, dsts, blockN);
   }
 
   // Even if there was nothing for this GPU to do, it must participate in a barrier
diff --git a/src/collectives/device/common_kernel.h b/src/collectives/device/common_kernel.h
index 9351c19a76..525e3a4cda 100644
--- a/src/collectives/device/common_kernel.h
+++ b/src/collectives/device/common_kernel.h
@@ -503,7 +503,7 @@ inline __device__ void Store128(Pack128* p, Pack128& v) {
 
 template<class FUNC, typename T, int UNROLL, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
 __device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const int t,
-    FUNC fn, bool preOpSrc0, bool postOp, int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Nelem
+    FUNC fn, int const numPreOpSrcs, bool postOp, int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Nelem
   ) {
   const int inc = nw * UNROLL * WARP_SIZE;
   int offset = w * UNROLL * WARP_SIZE + t;
@@ -517,7 +517,7 @@ __device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const
     T vals[UNROLL];
     // Load and reduce
     for (int u = 0; u < UNROLL; ++u) vals[u] = vFetch(srcs[0]+u*WARP_SIZE);
-    if (preOpSrc0) {
+    if (numPreOpSrcs) {
       for (int u = 0; u < UNROLL; ++u) vals[u] = FuncTraits<FUNC>().preOp(fn, vals[u]);
     }
 
@@ -525,6 +525,9 @@ __device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const
     for (int i=1; i<MINSRCS; i++) {
       T vals2[UNROLL];
       for (int u = 0; u < UNROLL; ++u) vals2[u] = vFetch(srcs[i]+u*WARP_SIZE);
+      if (i < numPreOpSrcs) {
+        for (int u = 0; u < UNROLL; ++u) vals2[u] = FuncTraits<FUNC>().preOp(fn, vals2[u]);
+      }
       for (int u = 0; u < UNROLL; ++u) vals[u] = fn(vals[u], vals2[u]);
     }
     #pragma unroll
@@ -560,7 +563,7 @@ __device__ __forceinline__ void ReduceCopyMulti(const int w, const int nw, const
 
 template<class FUNC, typename T, int UNROLL, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
 __device__ __forceinline__ void ReduceCopy128bMulti(const int w, const int nw, const int t,
-    FUNC fn, bool preOpSrc0, bool postOp, int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Npack
+    FUNC fn, int numPreOpSrcs, bool postOp, int nsrcs, const T** s, int ndsts, T** d, const int elemOffset, const int Npack
   ) {
   const int inc = nw * UNROLL * WARP_SIZE;
   int offset = w * UNROLL * WARP_SIZE + t;
@@ -574,7 +577,7 @@ __device__ __forceinline__ void ReduceCopy128bMulti(const int w, const int nw, c
     Pack128 vals[UNROLL];
     // Load and reduce
     for (int u = 0; u < UNROLL; ++u) Fetch128(vals[u], srcs[0]+u*WARP_SIZE);
-    if (preOpSrc0) {
+    if (numPreOpSrcs) {
       for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>().preOp(fn, vals[u]);
     }
 
@@ -582,6 +585,9 @@ __device__ __forceinline__ void ReduceCopy128bMulti(const int w, const int nw, c
     for (int i=1; i<MINSRCS; i++) {
       Pack128 vals2[UNROLL];
       for (int u = 0; u < UNROLL; ++u) Fetch128(vals2[u], srcs[i]+u*WARP_SIZE);
+      if (i < numPreOpSrcs) {
+        for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>().preOp(fn, vals2[u]);
+      }
       for (int u = 0; u < UNROLL; ++u) MULTI128<FUNC, T>()(fn, vals[u], vals2[u]);
     }
     #pragma unroll
@@ -623,7 +629,7 @@ __device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(uint32_t); }
 
 template<int UNROLL, class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
 __device__ __forceinline__ void ReduceOrCopyMulti(
-    const int tid, const int nthreads, FUNC fn, bool preOpSrc0, bool postOp, int nsrcs, const T** srcs, int ndsts, T** dsts, int N
+    const int tid, const int nthreads, FUNC fn, int numPreOpSrcs, bool postOp, int nsrcs, const T** srcs, int ndsts, T** dsts, int N
   ) {
   int Nrem = N;
   if (Nrem <= 0) return;
@@ -651,7 +657,7 @@ __device__ __forceinline__ void ReduceOrCopyMulti(
     int Nelem = Npack * PACKELEMS;
 
     ReduceCopy128bMulti<FUNC, T, AUTOUNROLL, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
-      (w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Npack);
+      (w, nw, t, fn, numPreOpSrcs, postOp, nsrcs, srcs, ndsts, dsts, offset, Npack);
 
     Nrem -= Nelem;
     if (Nrem == 0) return;
@@ -662,7 +668,7 @@ __device__ __forceinline__ void ReduceOrCopyMulti(
     Nelem = Npack * PACKELEMS;
 
     ReduceCopy128bMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
-      (w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Npack);
+      (w, nw, t, fn, numPreOpSrcs, postOp, nsrcs, srcs, ndsts, dsts, offset, Npack);
 
     Nrem -= Nelem;
     if (Nrem == 0) return;
@@ -673,7 +679,7 @@ __device__ __forceinline__ void ReduceOrCopyMulti(
   int Nelem = (Nrem / (AUTOUNROLL*PACKELEMS/2*WARP_SIZE)) * (AUTOUNROLL*PACKELEMS/2*WARP_SIZE); // round down
 
   ReduceCopyMulti<FUNC, T, AUTOUNROLL*PACKELEMS/2, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
-    (w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Nelem);
+    (w, nw, t, fn, numPreOpSrcs, postOp, nsrcs, srcs, ndsts, dsts, offset, Nelem);
 
   Nrem -= Nelem;
   if (Nrem == 0) return;
@@ -681,7 +687,7 @@ __device__ __forceinline__ void ReduceOrCopyMulti(
 
   // no unroll, by type. Should finish what's remaining.
   ReduceCopyMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>
-    (w, nw, t, fn, preOpSrc0, postOp, nsrcs, srcs, ndsts, dsts, offset, Nrem);
+    (w, nw, t, fn, numPreOpSrcs, postOp, nsrcs, srcs, ndsts, dsts, offset, Nrem);
 }
 
 #endif // COMMON_KERNEL_H_
diff --git a/src/collectives/device/prims_simple.h b/src/collectives/device/prims_simple.h
index 12ba088fea..9db8013542 100644
--- a/src/collectives/device/prims_simple.h
+++ b/src/collectives/device/prims_simple.h
@@ -175,14 +175,14 @@ class Primitives<
           if (Send) {
             // (1-Send) is only there to avoid compilation errors in case MaxSend=0 (and Send=0).
             ReduceOrCopyMulti<Unroll, RedOp, T, 1, 1, 1, (1-Send)+MaxSend>
-              (tid, nworkers, redOp, false, false,
+              (tid, nworkers, redOp, 0, false,
                1, (T const**)ncclShmem->groups[group].srcs,
                fan.nsend(), (T**)ncclShmem->groups[group].dsts+1,
                sliceSize);
           }
         } else {
           ReduceOrCopyMulti<Unroll, RedOp, T, Recv+Src, Recv*MaxRecv+Src, Send+Dst, Send*MaxSend+Dst>
-            (tid, nworkers, redOp, SrcBuf==Input, postOp,
+            (tid, nworkers, redOp, SrcBuf==Input ? 1 : 0, postOp,
              Recv*fan.nrecv()+Src, (T const**)ncclShmem->groups[group].srcs,
              Send*fan.nsend()+Dst, (T**)ncclShmem->groups[group].dsts,
              sliceSize);
@@ -241,7 +241,7 @@ class Primitives<
             if (skip >= 0 && i >= skip) peerOffset += peerElem;
             const T* src0 = (T*)ncclShmem->groups[group].srcs[0] + peerOffset;
             int realPeerSize = min(realSize, totalElem-peerOffset);
-            if (realPeerSize > 0) ReduceOrCopyMulti<Unroll, RedOp, T, 1, 1, 1, 1>(tid, nworkers, redOp, true, false, 1, &src0, 1, (T**)ncclShmem->groups[group].dsts+i, realPeerSize);
+            if (realPeerSize > 0) ReduceOrCopyMulti<Unroll, RedOp, T, 1, 1, 1, 1>(tid, nworkers, redOp, 1, false, 1, &src0, 1, (T**)ncclShmem->groups[group].dsts+i, realPeerSize);
           }
         } else if (Recv) {
           #pragma unroll 1
@@ -251,7 +251,7 @@ class Primitives<
             if (skip >= 0 && i >= skip) peerOffset += peerElem;
             T* dst0 = (T*)ncclShmem->groups[group].dsts[0] + peerOffset;
             int realPeerSize = min(realSize, totalElem-peerOffset);
-            if (realPeerSize > 0) ReduceOrCopyMulti<Unroll, RedOp, T, 1, 1, 1, 1>(tid, nworkers, redOp, false, postOp, 1, (T const**)ncclShmem->groups[group].srcs+i, 1, &dst0, realPeerSize);
+            if (realPeerSize > 0) ReduceOrCopyMulti<Unroll, RedOp, T, 1, 1, 1, 1>(tid, nworkers, redOp, 0, postOp, 1, (T const**)ncclShmem->groups[group].srcs+i, 1, &dst0, realPeerSize);
           }
         }
       }
diff --git a/src/collectives/device/sendrecv.h b/src/collectives/device/sendrecv.h
index 7e8b6eb5a1..f717297df4 100644
--- a/src/collectives/device/sendrecv.h
+++ b/src/collectives/device/sendrecv.h
@@ -45,7 +45,7 @@ struct RunWork<ncclFuncSendRecv, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_SIMPLE> {
             for (size_t offset=0; offset<sendCount; offset += blockSize) {
               size_t remaining = sendCount - offset;
               if (remaining < blockSize) blockSize = remaining;
-              ReduceOrCopyMulti<COLL_UNROLL, RedOp, T, 1, 1, 1, 1>(tid, nThreadsSegment, RedOp(), false, false, 1, &sendbuff, 1, &recvbuff, blockSize);
+              ReduceOrCopyMulti<COLL_UNROLL, RedOp, T, 1, 1, 1, 1>(tid, nThreadsSegment, RedOp(), 0, false, 1, &sendbuff, 1, &recvbuff, blockSize);
               sendbuff += blockSize;
               recvbuff += blockSize;
             }