From e6b6a3752842d00fc6061f8dc8ae218b7e7bb552 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Tue, 4 Feb 2025 13:14:43 -0800
Subject: [PATCH] Revert "Remove unused code path (#1527)" (#1530)

This reverts commit a7d9bfda6e63c4d86add8c81951820329d3be7e2.

[ROCm/rccl commit: d00e903d7224d0a3d16efcec594c65791082e3ea]
---
 projects/rccl/src/device/prims_simple.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/rccl/src/device/prims_simple.h b/projects/rccl/src/device/prims_simple.h
index 11251684ee..1887ea154b 100644
--- a/projects/rccl/src/device/prims_simple.h
+++ b/projects/rccl/src/device/prims_simple.h
@@ -369,7 +369,7 @@ private:
     // worker perf is the limiter, perf-wise this loop is effectively unentered,
     // hence just a single branch insn.
     #pragma unroll 1
-    while (false) {
+    while (slice < SlicePerChunk) {
       sliceSize = sliceSize < nelem-offset ? sliceSize : nelem-offset;
       { // Only workers could have Wait roles so we know the slice must be empty
         // since we've exited the loop above.
@@ -942,4 +942,4 @@ private:
   __device__ __forceinline__ void localCopy(T* srcs, T* dsts, int eltN) {
     return mscclGenericOp<0,1,0,0>(&srcs, 1, &dsts, 1, eltN);
   }
-};
+};
\ No newline at end of file