From d1e9058c8e0a0f25de7536282bfa0bb6e833670c 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> --- CMakeLists.txt | 2 +- include/ck/ck.hpp | 14 ++++++++------ .../impl/device_batched_gemm_multiple_d_dl.hpp | 4 ++-- .../gpu/device/impl/device_gemm_multiple_d_dl.hpp | 2 +- .../impl/device_grouped_conv_bwd_weight_dl.hpp | 4 ++-- ...ouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp | 4 ++-- .../device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp | 4 ++-- .../impl/device_grouped_gemm_multiple_d_dl.hpp | 4 ++-- .../gpu/grid/gridwise_tensor_rearrange.hpp | 4 ++-- include/ck_tile/core/config.hpp | 6 +++++- profiler/src/CMakeLists.txt | 9 +++++++-- 11 files changed, 34 insertions(+), 23 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index b0fc725236..e0ff729df8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -101,7 +101,7 @@ add_compile_options(-Wno-unique-object-duplication) # 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) +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 3c1373a387..40d1abbc28 100644 --- a/include/ck/ck.hpp +++ b/include/ck/ck.hpp @@ -56,7 +56,8 @@ #if defined(__gfx942__) || defined(__gfx950__) #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__) || \ @@ -78,20 +79,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 1b487502f4..203b2a6168 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,8 +71,8 @@ __global__ void const Block2CTileMap block_2_ctile_map) { #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx908__) || \ - defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx103__) || defined(__gfx11__) || \ - defined(__gfx12__)) + defined(__gfx90a__) || defined(__gfx94__) || 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 b9467ac194..9be0916d68 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 @@ -51,7 +51,7 @@ __global__ void const Block2CTileMap block_2_ctile_map) { #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx9__) || \ - defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) + 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 0b3f1a0255..271bbd6fc7 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,8 +49,8 @@ __global__ void const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx103__) || \ - defined(__gfx90a__) || defined(__gfx908__) || defined(__gfx94__) || defined(__gfx11__) || \ - defined(__gfx12__)) + defined(__gfx90a__) || defined(__gfx908__) || 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 3e14f66a09..d36f223773 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 @@ -91,8 +91,8 @@ __global__ void const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx103__) || \ - defined(__gfx90a__) || defined(__gfx908__) || defined(__gfx94__) || defined(__gfx11__) || \ - defined(__gfx12__)) + defined(__gfx90a__) || defined(__gfx908__) || 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 50e171e503..1eb0d7255e 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,8 +106,8 @@ __global__ void const Block2CTileMap block_2_ctile_map, const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) { -#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx103__) || \ - defined(__gfx11__) || defined(__gfx12__)) +#if(!defined(__HIP_DEVICE_COMPILE__) || 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 10d8a4a44d..f22bf7a635 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 @@ -41,8 +41,8 @@ __global__ void const CDEElementwiseOperation cde_element_op) { #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx908__) || \ - defined(__gfx90a__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx94__) || \ - defined(__gfx12__)) + 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 ddf0b4a58d..2d9a7abce2 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp @@ -36,8 +36,8 @@ __global__ void const ComputePtrOffsetOfStridedBatch compute_ptr_offset_of_batch) { #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx908__) || \ - defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx103__) || defined(__gfx11__) || \ - defined(__gfx12__)) + defined(__gfx90a__) || defined(__gfx94__) || 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 1ecc28fbeb..aa8b72b9e9 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -9,6 +9,10 @@ #if defined(__gfx942__) || defined(__gfx950__) #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__) @@ -200,7 +204,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 fef09315d5..69bc78bd0d 100644 --- a/profiler/src/CMakeLists.txt +++ b/profiler/src/CMakeLists.txt @@ -223,11 +223,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}")