diff --git a/projects/rccl/src/collectives/device/prims_ll.h b/projects/rccl/src/collectives/device/prims_ll.h index 0c96ca9022..388cec161d 100644 --- a/projects/rccl/src/collectives/device/prims_ll.h +++ b/projects/rccl/src/collectives/device/prims_ll.h @@ -183,7 +183,7 @@ private: template __device__ void readLLBeginAll(int offset, ncclLLFifoLine(&line)[MaxRecv]) { - #pragma unroll 1 + #pragma unroll for (int i=BeginIx; i < MaxRecv; i++) { if (i < fan.nrecv()) { union ncclLLFifoLine* src = recvPtr(i) + offset; @@ -412,7 +412,7 @@ private: } if (RECV) { data = !SRC ? peerData : MULTI()(redOp, peerData, data); - #pragma unroll 1 + #pragma unroll for (int i=1; i < MaxRecv && i < fan.nrecv(); i++) { peerData = readLLFinish(offset, line, i); data = MULTI()(redOp, peerData, data); diff --git a/projects/rccl/src/collectives/device/prims_simple.h b/projects/rccl/src/collectives/device/prims_simple.h index aafb0ec9b5..b4c11fa9a7 100644 --- a/projects/rccl/src/collectives/device/prims_simple.h +++ b/projects/rccl/src/collectives/device/prims_simple.h @@ -200,7 +200,6 @@ private: // barrier(); // post(); // } // Since we no longer unroll, new branch added here - #pragma unroll 1 do { sliceSize = sliceSize < nelem-offset ? sliceSize : nelem-offset; if (Src && (flags & (SrcBuf==Input ? RoleInput : RoleOutput))) @@ -339,7 +338,6 @@ private: // slices are all empty. Since empty slices are the uncommon case, and // worker perf is the limiter, perf-wise this loop is effectively unentered, // hence just a single branch insn. - #pragma unroll 1 while (slice < SlicePerChunk) { sliceSize = sliceSize < nelem-offset ? sliceSize : nelem-offset; { // Only workers could have Wait roles so we know the slice must be empty @@ -371,7 +369,6 @@ private: int sliceSize = stepSize*StepPerSlice; int dataSize = max(DIVUP(peerElem, 16*SlicePerChunk)*16, sliceSize/32); // per-peer slice size - #pragma unroll 1 for (int slice=0; slice(0, inpIx, offset, realSize); subBarrier(); - #pragma unroll 1 // Loop over peers for (int j=0; j()(vals[u], vals2[u]); } - #pragma unroll 1 for (int i=MINSRCS; i(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Nrem); } -#endif \ No newline at end of file +#endif