AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Disable LL128 for gfx90a (#1054)

Telusuri Sumber 
[ROCm/rccl commit: 1ac800f9fc]
This commit is contained in:
Bertan Dogancay
2024-01-18 18:34:19 -07:00
committed by GitHub
orang tua 4be87126fe
melakukan c91be550bc
1 mengubah file dengan 36 tambahan dan 2 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/cmake/Generator.cmake
+36 -2
Melihat File
@@ -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})