diff --git a/projects/rccl/cmake/Generator.cmake b/projects/rccl/cmake/Generator.cmake
index 7871d01337..7fb41209d9 100644
--- a/projects/rccl/cmake/Generator.cmake
+++ b/projects/rccl/cmake/Generator.cmake
@@ -154,11 +154,18 @@ function(gen_device_table)
 
   ## Declaration of device functions
   foreach(func IN LISTS FUNC_LIST)
+    string(FIND "${func}" "LL128" IS_LL128)
+    if(NOT IS_LL128 EQUAL -1)
+      file(APPEND ${DEVICE_TABLE_FILE} "#if defined(__gfx90a__)\n")
+    endif()
     if(ENABLE_IFC)
       file(APPEND ${DEVICE_TABLE_FILE} "__device__ void ${func}();\n")
     else()
       file(APPEND ${DEVICE_TABLE_FILE} "__device__ __attribute__((noinline)) void ${func}();\n")
     endif()
+    if(NOT IS_LL128 EQUAL -1)
+      file(APPEND ${DEVICE_TABLE_FILE} "#endif\n")
+    endif()
   endforeach()
   file(APPEND ${DEVICE_TABLE_FILE} "\n")
 
@@ -166,7 +173,17 @@ function(gen_device_table)
     ## Undirect function call
     file(APPEND ${DEVICE_TABLE_FILE} "__device__ ncclKernelFunc_t const ncclFuncs[] = {\n")
     foreach(func ${FUNC_LIST})
-      file(APPEND ${DEVICE_TABLE_FILE} "  ${func},\n")
+      string(FIND "${func}" "LL128" IS_LL128)
+      if(NOT IS_LL128 EQUAL -1)
+        file(APPEND ${DEVICE_TABLE_FILE} "#if defined(__gfx90a__)\n")
+        file(APPEND ${DEVICE_TABLE_FILE} "  ${func},\n")
+        file(APPEND ${DEVICE_TABLE_FILE} "#else\n")
+        string(REPLACE "LL128" "LL" func "${func}")
+        file(APPEND ${DEVICE_TABLE_FILE} "  ${func},\n")
+        file(APPEND ${DEVICE_TABLE_FILE} "#endif\n")
+      else()
+        file(APPEND ${DEVICE_TABLE_FILE} "  ${func},\n")
+      endif()
     endforeach()
     ## Add OneRankReduce functions at the end
     foreach(type IN LISTS ALL_TYPES)
@@ -178,7 +195,17 @@ function(gen_device_table)
     file(APPEND ${DEVICE_TABLE_FILE} "__device__ void NCCL_CALL_FUNCTIONS(unsigned short funcIndex) noexcept {\n  switch(funcIndex) {\n")
     set(index 0)
     foreach(func IN LISTS FUNC_LIST)
-      file(APPEND ${DEVICE_TABLE_FILE} "    case ${index}:\n      ${func}();\n      break;\n")
+      string(FIND "${func}" "LL128" IS_LL128)
+      if(NOT IS_LL128 EQUAL -1)
+        file(APPEND ${DEVICE_TABLE_FILE} "#if defined(__gfx90a__)\n")
+        file(APPEND ${DEVICE_TABLE_FILE} "    case ${index}:\n      ${func}();\n      break;\n")
+        file(APPEND ${DEVICE_TABLE_FILE} "#else\n")
+        string(REPLACE "LL128" "LL" func "${func}")
+        file(APPEND ${DEVICE_TABLE_FILE} "    case ${index}:\n      ${func}();\n      break;\n")
+        file(APPEND ${DEVICE_TABLE_FILE} "#endif\n")
+      else()
+        file(APPEND ${DEVICE_TABLE_FILE} "    case ${index}:\n      ${func}();\n      break;\n")
+      endif()
       math(EXPR index "${index} + 1")
     endforeach()
     ## Add OneRankReduce functions at the end
@@ -334,9 +361,16 @@ function(gen_collectives)
 
     ## Construct the file
     file(WRITE ${FILE_PATH} "#include \"${COLL_LOWER}.h\"\n#include \"common.h\"\n#include \"collectives.h\"\n")
+    string(FIND "${list_name}" "LL128" IS_LL128)
+    if(NOT IS_LL128 EQUAL -1)
+      file(APPEND ${FILE_PATH} "#if defined(__gfx90a__)\n")
+    endif()
     foreach(IMPL IN LISTS IMPL_LIST)
       file(APPEND ${FILE_PATH} "${IMPL}")
     endforeach()
+    if(NOT IS_LL128 EQUAL -1)
+      file(APPEND ${FILE_PATH} "#endif\n")
+    endif()
 
     ## Append the file to HIP sources list which will be added to source list
     list(APPEND HIP_SOURCES ${FILE_PATH})