From b6f47bad7c63040a5a216c94e02de769c3d3af9f Mon Sep 17 00:00:00 2001
From: Nusrat Islam <Nusrat.Islam@amd.com>
Date: Thu, 2 Nov 2023 21:41:39 -0500
Subject: [PATCH] msccl: remove cases from numReduction switch statement

[ROCm/rccl commit: f545b94d4b200a50e7d386f205512b640ff6a3fe]
---
 .../collectives/device/msccl_kernel_impl.h    | 20 ++++---------------
 1 file changed, 4 insertions(+), 16 deletions(-)

diff --git a/projects/rccl/src/collectives/device/msccl_kernel_impl.h b/projects/rccl/src/collectives/device/msccl_kernel_impl.h
index 9699b1d029..90b9051dc9 100644
--- a/projects/rccl/src/collectives/device/msccl_kernel_impl.h
+++ b/projects/rccl/src/collectives/device/msccl_kernel_impl.h
@@ -352,22 +352,16 @@ __device__ __forceinline__ void mscclRunInterpreter(
               T o = load(dstIndex);
               ssize_t srcBaseOffset = gridOffset + (ssize_t)c * sizePerMscclChunk + tid;
               switch (numReductions) {
-                case 1:
-                  #pragma unroll
-                  MSCCL_REDUCE_UNROLL_LOOP_A(1);
-                  break;
-                case 3:
-                  #pragma unroll
-                  MSCCL_REDUCE_UNROLL_LOOP_A(3);
-                  break;
                 case 7:
                   #pragma unroll
                   MSCCL_REDUCE_UNROLL_LOOP_A(7);
                   break;
+#if defined(__gfx90a__)
                 case 15:
                   #pragma unroll
                   MSCCL_REDUCE_UNROLL_LOOP_A(15);
                   break;
+#endif
                 default:
                   MSCCL_REDUCE_UNROLL_LOOP_A(numReductions);
                   break;
@@ -388,22 +382,16 @@ __device__ __forceinline__ void mscclRunInterpreter(
             T* dst = dstPointer + dstOffset;
             ssize_t srcBaseOffset = gridOffset + (ssize_t)c * sizePerMscclChunk;
             switch (numReductions) {
-              case 1:
-                #pragma unroll
-                MSCCL_REDUCE_UNROLL_LOOP_B(1);
-                break;
-              case 3:
-                #pragma unroll
-                MSCCL_REDUCE_UNROLL_LOOP_B(3);
-                break;
               case 7:
                 #pragma unroll
                 MSCCL_REDUCE_UNROLL_LOOP_B(7);
                 break;
+#if defined(__gfx90a__)
               case 15:
                 #pragma unroll
                 MSCCL_REDUCE_UNROLL_LOOP_B(15);
                 break;
+#endif
               default:
                 MSCCL_REDUCE_UNROLL_LOOP_B(numReductions);
                 break;