From 420f8af6a0f167a4b062b45fd3e5ef607c96b965 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Wed, 26 Jul 2023 11:44:15 -0700
Subject: [PATCH] Enable Ll128 on gfx90a (#823)

---
 src/collectives/device/common.h | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/src/collectives/device/common.h b/src/collectives/device/common.h
index 18cf382652..82d86cb542 100644
--- a/src/collectives/device/common.h
+++ b/src/collectives/device/common.h
@@ -28,7 +28,7 @@
   { __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST); }
 #endif
 
-#ifdef ENABLE_LL128
+#ifdef __gfx90a__
 #define NCCL_FUNC5(func, algo, devredop, type, nullify) \
   MACRO_IF(nullify, nullptr, NCCL_FUNC_NAME(func, algo, LL,     devredop, type)), \
   MACRO_IF(nullify, nullptr, NCCL_FUNC_NAME(func, algo, LL128,  devredop, type)), \
@@ -549,7 +549,7 @@ __forceinline__ __device__ void ncclKernel(
 #ifdef USE_INDIRECT_FUNCTION_CALL
       ncclFuncs[ncclShmem.work.header.funcIndex]();
 #else
-#ifdef ENABLE_LL128
+#ifdef __gfx90a__
       NCCL_CALL_FUNCTIONS<1>(ncclShmem.work.header.funcIndex);
 #else
       NCCL_CALL_FUNCTIONS<0>(ncclShmem.work.header.funcIndex);
@@ -618,7 +618,7 @@ __device__  __attribute__((noinline)) void NCCL_FUNC_NAME(func, algo, proto, dev
 #endif
 
 // Only generate inline kernels for LL
-#ifdef ENABLE_LL128
+#ifdef __gfx90a__
 #define IMPL_COLL4(func, algo, devredop, type) \
   IMPL_COLL_FUNC(func, algo, LL,     devredop, type) \
   IMPL_COLL_FUNC(func, algo, LL128,  devredop, type) \