From d80f38f77f6dd5087c102488e5ff9378e0777748 Mon Sep 17 00:00:00 2001 From: Gavin Zhao Date: Thu, 20 Nov 2025 13:45:57 -0500 Subject: [PATCH] Add support for RDNA1 GPUs (#3220) * Allow compilation for RDNA1 (__gfx101__) Signed-off-by: Gavin Zhao * More RDNA1 changes Signed-off-by: Gavin Zhao * Even more RDNA1 changes Signed-off-by: Gavin Zhao * cmake: skip build quantization for unsupported arches * add gfx10-1-generic support as well * add gfx1013 and complete gfx10-1-generic * fix clang format * enable DL kernels on gfx101x --------- Signed-off-by: Gavin Zhao Co-authored-by: illsilin_amdeng Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com> [ROCm/composable_kernel commit: 07314ac54374d6e1e5fb2785fa8d58ceac21d81c] --- CMakeLists.txt | 2 +- include/ck/ck.hpp | 14 ++++++++------ .../impl/device_batched_gemm_multiple_d_dl.hpp | 2 +- .../gpu/device/impl/device_gemm_multiple_d_dl.hpp | 4 ++-- .../impl/device_grouped_conv_bwd_weight_dl.hpp | 2 +- ...ouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp | 2 +- .../device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp | 3 ++- .../impl/device_grouped_gemm_multiple_d_dl.hpp | 4 ++-- .../gpu/grid/gridwise_tensor_rearrange.hpp | 2 +- include/ck_tile/core/config.hpp | 6 +++++- profiler/src/CMakeLists.txt | 9 +++++++-- test/quantization/gemm/CMakeLists.txt | 13 +++++++++---- 12 files changed, 40 insertions(+), 23 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index b1532f2cc8..45db703b82 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -122,7 +122,7 @@ add_compile_options( # Recent change in compiler makes this warning ON by default, which led to compile errors. add_compile_options(-Wno-nrvo) -if(NOT DISABLE_DL_KERNELS AND GPU_TARGETS MATCHES "gfx103|gfx10-3-generic") +if(NOT DISABLE_DL_KERNELS AND GPU_TARGETS MATCHES "gfx101|gfx103|gfx10-1|gfx10-3") add_definitions(-DDL_KERNELS) set(DL_KERNELS "ON") set(CK_ENABLE_DL_KERNELS "ON") diff --git a/include/ck/ck.hpp b/include/ck/ck.hpp index e4ad477dde..9debdc12b2 100644 --- a/include/ck/ck.hpp +++ b/include/ck/ck.hpp @@ -58,7 +58,8 @@ #if defined(__gfx942__) || defined(__gfx950__) || defined(__gfx9_4_generic__) #define __gfx94__ #endif -#if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) +#if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || \ + defined(__gfx1013__) || defined(__gfx10_1_generic__) #define __gfx101__ #endif #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \ @@ -80,20 +81,21 @@ #define CK_BUFFER_RESOURCE_3RD_DWORD -1 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__) #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000 -#elif defined(__gfx103__) +#elif defined(__gfx101__) || defined(__gfx103__) #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000 #elif defined(__gfx11__) || defined(__gfx12__) #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000 #endif // FMA instruction -#ifndef __HIP_DEVICE_COMPILE__ // for host code, define nothing -#elif defined(__gfx803__) || defined(__gfx900__) // for GPU code -#define CK_USE_AMD_V_MAC_F32 -#elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) // for GPU code +#ifndef __HIP_DEVICE_COMPILE__ // for host code, define nothing +#elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) || defined(__gfx1011__) || \ + defined(__gfx1012__) // for GPU code #define CK_USE_AMD_V_FMAC_F32 #define CK_USE_AMD_V_DOT2_F32_F16 #define CK_USE_AMD_V_DOT4_I32_I8 +#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx101__) +#define CK_USE_AMD_V_MAC_F32 #elif defined(__gfx11__) || defined(__gfx12__) #define CK_USE_AMD_V_FMAC_F32 #define CK_USE_AMD_V_DOT2_F32_F16 diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_dl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_dl.hpp index 00004cc9a9..944c93acdc 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_dl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_dl.hpp @@ -71,7 +71,7 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const Block2CTileMap block_2_ctile_map) { #if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || \ - defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) + defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) const index_t num_blocks_per_batch = __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count); diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_dl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_dl.hpp index 13de33f80f..11d1a74819 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_dl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_dl.hpp @@ -50,8 +50,8 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const CGridDesc_M0_M10_M11_N0_N10_N11 e_grid_desc_m0_m10_m11_n0_n10_n11, const Block2CTileMap block_2_ctile_map) { -#if(defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) || defined(__gfx11__) || \ - defined(__gfx12__)) +#if(defined(__gfx906__) || defined(__gfx9__) || defined(__gfx101__) || defined(__gfx103__) || \ + defined(__gfx11__) || defined(__gfx12__)) constexpr index_t shared_block_size = GridwiseGemm::GetSharedMemoryNumberOfByte() / sizeof(ABDataType); diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp index 4e0b05817f..2152a72105 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp @@ -49,7 +49,7 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { #if(defined(__gfx906__) || defined(__gfx103__) || defined(__gfx90a__) || defined(__gfx908__) || \ - defined(__gfx94__) || defined(__gfx11__) || defined(__gfx12__)) + defined(__gfx94__) || defined(__gfx101__) || defined(__gfx11__) || defined(__gfx12__)) const index_t num_blocks_per_batch = __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count); const index_t g_idx = __builtin_amdgcn_readfirstlane(get_block_1d_id() / num_blocks_per_batch); diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp index 572c52b8fd..347ea25e62 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp @@ -95,7 +95,7 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { #if(defined(__gfx906__) || defined(__gfx103__) || defined(__gfx90a__) || defined(__gfx908__) || \ - defined(__gfx94__) || defined(__gfx11__) || defined(__gfx12__)) + defined(__gfx94__) || defined(__gfx101__) || defined(__gfx11__) || defined(__gfx12__)) // offset base pointer for each work-group const index_t num_blocks_per_batch = __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count); diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp index 9278ceeb58..b1e8f74ce9 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp @@ -106,7 +106,8 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const Block2CTileMap block_2_ctile_map, const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { -#if(defined(__gfx906__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) +#if(defined(__gfx906__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || \ + defined(__gfx12__)) // offset base pointer for each work-group const index_t num_blocks_per_batch = __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count); diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp index da4c416794..b7b5b622ff 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp @@ -39,8 +39,8 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const BElementwiseOperation b_element_op, const CDEElementwiseOperation cde_element_op) { -#if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx103__) || \ - defined(__gfx11__) || defined(__gfx94__) || defined(__gfx12__)) +#if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx101__) || \ + defined(__gfx103__) || defined(__gfx11__) || defined(__gfx94__) || defined(__gfx12__)) __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()]; const index_t block_id = get_block_1d_id(); diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp index 295a77ca34..9761cc6a68 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp @@ -36,7 +36,7 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU) const ComputePtrOffsetOfStridedBatch compute_ptr_offset_of_batch) { #if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || \ - defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) + defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) GridwiseTensorRearrangeKernel::Run(in_grid_desc, p_in_global, out_grid_desc, diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index 76a1c03269..b01f9dedef 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -10,6 +10,10 @@ #if defined(__gfx942__) || defined(__gfx950__) || defined(__gfx9_4_generic__) #define __gfx94__ #endif +#if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || \ + defined(__gfx1013__) || defined(__gfx10_1_generic__) +#define __gfx101__ +#endif #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \ defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \ defined(__gfx10_3_generic__) @@ -211,7 +215,7 @@ #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || \ defined(__gfx9__) // for GPU code #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x00020000 -#elif defined(__gfx103__) // for GPU code +#elif defined(__gfx101__) || defined(__gfx103__) // for GPU code #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31014000 #elif defined(__gfx11__) || defined(__gfx12__) // for GPU code #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31004000 diff --git a/profiler/src/CMakeLists.txt b/profiler/src/CMakeLists.txt index c22867fbed..b9f82af29d 100644 --- a/profiler/src/CMakeLists.txt +++ b/profiler/src/CMakeLists.txt @@ -264,11 +264,16 @@ endif() set(PROFILER_LIBS utility getopt::getopt) foreach(LIB ${DEVICE_INSTANCES}) string(REGEX REPLACE "device_(.+)_instance" "\\1" INSTANCE_NAME ${LIB}) - if (INSTANCE_NAME STREQUAL "") + if (INSTANCE_NAME STREQUAL "") message(FATAL_ERROR "Unexpected kernel instance name: ${LIB}") endif() if("${INSTANCE_NAME}" MATCHES "${CK_PROFILER_INSTANCE_FILTER}") - list(APPEND PROFILER_LIBS ${LIB}) + # Only link if the target was actually created + if(TARGET ${LIB}) + list(APPEND PROFILER_LIBS ${LIB}) + else() + message(VERBOSE "Skipping ${LIB} - no instances built for current GPU targets") + endif() endif() endforeach() message(VERBOSE "ckProfiler libs: ${PROFILER_LIBS}") diff --git a/test/quantization/gemm/CMakeLists.txt b/test/quantization/gemm/CMakeLists.txt index 630e6e09c9..0eb08f9a5b 100644 --- a/test/quantization/gemm/CMakeLists.txt +++ b/test/quantization/gemm/CMakeLists.txt @@ -1,9 +1,14 @@ add_custom_target(test_gemm_quantization_targets) -add_gtest_executable(test_gemm_quantization test_gemm_quantization.cpp) -if(result EQUAL 0) - target_link_libraries(test_gemm_quantization PRIVATE utility device_quantization_instance) - add_dependencies(test_gemm_quantization_targets test_gemm_quantization) +# Only build test if the quantization instance library exists +if(TARGET device_quantization_instance) + add_gtest_executable(test_gemm_quantization test_gemm_quantization.cpp) + if(result EQUAL 0) + target_link_libraries(test_gemm_quantization PRIVATE utility device_quantization_instance) + add_dependencies(test_gemm_quantization_targets test_gemm_quantization) + endif() +else() + message(VERBOSE "Skipping test_gemm_quantization - device_quantization_instance not built for current GPU targets") endif() add_dependencies(test_quantization test_gemm_quantization_targets)