Skip to content

Commit 108171a

Browse files
slojosic-amdillsilin
slojosic-amd
and
illsilin
authored
enable gfx115x support (#2065) (#2552)
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
1 parent a8c5bd9 commit 108171a

6 files changed

Lines changed: 20 additions & 16 deletions

File tree

‎example/CMakeLists.txt‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -107,11 +107,11 @@ function(add_example_executable EXAMPLE_NAME FILE_NAME)
107107
#only continue if there are some source files left on the list
108108
if(FILE_NAME)
109109
if(FILE_NAME MATCHES "_xdl")
110-
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
110+
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
111111
elseif(FILE_NAME MATCHES "_wmma")
112112
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx950)
113113
elseif(FILE_NAME MATCHES "_mx") #only build mx example for gfx950
114-
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
114+
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
115115
endif()
116116
set_source_files_properties(${FILE_NAME} PROPERTIES LANGUAGE HIP)
117117
add_executable(${EXAMPLE_NAME} ${FILE_NAME})
@@ -202,7 +202,7 @@ function(add_example_executable_no_testing EXAMPLE_NAME FILE_NAME)
202202
#only continue if there are some source files left on the list
203203
if(FILE_NAME)
204204
if(FILE_NAME MATCHES "_xdl")
205-
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
205+
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
206206
elseif(FILE_NAME MATCHES "_wmma")
207207
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx950)
208208
endif()

‎include/ck/ck.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -70,7 +70,8 @@ CK_DECLARE_ENV_VAR_BOOL(CK_LOGGING)
7070
#define __gfx103__
7171
#endif
7272
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
73-
defined(__gfx1103__) || defined(__gfx11_generic__)
73+
defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
74+
defined(__gfx1152__) || defined(__gfx11_generic__)
7475
#define __gfx11__
7576
#endif
7677
#if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)

‎include/ck/host_utility/device_prop.hpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -88,7 +88,9 @@ inline bool is_gfx103_supported()
8888
inline bool is_gfx11_supported()
8989
{
9090
return ck::get_device_name() == "gfx1100" || ck::get_device_name() == "gfx1101" ||
91-
ck::get_device_name() == "gfx1102" || ck::get_device_name() == "gfx1103";
91+
ck::get_device_name() == "gfx1102" || ck::get_device_name() == "gfx1103" ||
92+
ck::get_device_name() == "gfx1150" || ck::get_device_name() == "gfx1151" ||
93+
ck::get_device_name() == "gfx1152";
9294
}
9395

9496
inline bool is_gfx12_supported()

‎include/ck_tile/core/config.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,8 @@
1616
#define __gfx103__
1717
#endif
1818
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
19-
defined(__gfx1103__) || defined(__gfx11_generic__)
19+
defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
20+
defined(__gfx1152__) || defined(__gfx11_generic__)
2021
#define __gfx11__
2122
#endif
2223
#if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)

‎library/src/tensor_operation_instance/gpu/CMakeLists.txt‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -107,22 +107,22 @@ function(add_instance_library INSTANCE_NAME)
107107
foreach(source IN LISTS ARGN)
108108
set(INST_TARGETS ${SUPPORTED_GPU_TARGETS})
109109
if(source MATCHES "_xdl")
110-
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
110+
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
111111
elseif(source MATCHES "_wmma")
112112
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx950)
113113
elseif(source MATCHES "mha")
114-
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx908 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
114+
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx908 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
115115
endif()
116116
#only build the fp8 gemm instances for gfx908/90a if the build argument is set
117117
if(NOT CK_USE_FP8_ON_UNSUPPORTED_ARCH)
118118
if(source MATCHES "gemm_xdl_universal" AND source MATCHES "f8")
119-
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
119+
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
120120
endif()
121121
if(source MATCHES "gemm_multiply_multiply_f8")
122-
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
122+
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
123123
endif()
124124
if(source MATCHES "bached_gemm_multiply_multiply_f8")
125-
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
125+
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
126126
endif()
127127
endif()
128128
set(offload_targets)

‎test/CMakeLists.txt‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -98,11 +98,11 @@ function(add_test_executable TEST_NAME)
9898
#only continue if there are some source files left on the list
9999
if(ARGN)
100100
if(ARGN MATCHES "_xdl")
101-
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
101+
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
102102
elseif(ARGN MATCHES "_wmma")
103103
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx950)
104104
elseif(ARGN MATCHES "_smfmac")
105-
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx908 gfx90a gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
105+
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx908 gfx90a gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
106106
endif()
107107
set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
108108
add_executable(${TEST_NAME} ${ARGN})
@@ -194,13 +194,13 @@ function(add_gtest_executable TEST_NAME)
194194
#only continue if there are some source files left on the list
195195
if(ARGN)
196196
if(ARGN MATCHES "_xdl")
197-
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
197+
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
198198
elseif(ARGN MATCHES "_wmma")
199199
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx950)
200200
elseif(ARGN MATCHES "_smfmac")
201-
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx908 gfx90a gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
201+
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx908 gfx90a gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
202202
elseif(ARGN MATCHES "_mx") #only build mx example for gfx950
203-
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
203+
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic)
204204
endif()
205205
set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
206206
add_executable(${TEST_NAME} ${ARGN})

0 commit comments

Comments
 (0)