From 6b1d184e66e143a6fa0c1b7d049c373a082ecc4c Mon Sep 17 00:00:00 2001 From: xiangxli Date: Thu, 30 Apr 2026 12:33:00 +0000 Subject: [PATCH] tmp --- include/ck_tile/core/arch/arch.hpp | 5 +++++ include/ck_tile/core/config.hpp | 9 ++++++++- 2 files changed, 13 insertions(+), 1 deletion(-) diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp index 417ec12c8c..c1a45cf155 100644 --- a/include/ck_tile/core/arch/arch.hpp +++ b/include/ck_tile/core/arch/arch.hpp @@ -104,6 +104,7 @@ enum struct amdgcn_target_id GFX11_GENERIC = 0x11FF, GFX1200 = 0x1200, GFX1201 = 0x1201, + GFX1250 = 0x1250, GFX12_GENERIC = 0x12FF, HOST = 0x0000, }; @@ -301,6 +302,7 @@ constexpr auto get_compiler_target() MAP_COMPILER_STATE_TO_GFX11_TARGET(CK_TILE_ARCH_GFX11_GENERIC, GFX11_GENERIC); MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX1200, GFX1200); MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX1201, GFX1201); + MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX1250, GFX1250); MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX12_GENERIC, GFX12_GENERIC); // Return HOST by default @@ -369,6 +371,7 @@ CK_TILE_HOST auto hip_device_prop_gcn_arch_name_to_amdgcn_target_id(char const* MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx11_generic", GFX11_GENERIC); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1200", GFX1200); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1201", GFX1201); + MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1250", GFX1250); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx12_generic", GFX12_GENERIC); // Default case: return HOST target if no match is found @@ -626,6 +629,7 @@ CK_TILE_HOST_DEVICE constexpr auto get_compiler_target() MAP_COMPILER_STATE_TO_GFX11_TARGET(CK_TILE_ARCH_GFX11_GENERIC, GFX11_GENERIC); MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX1200, GFX1200); MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX1201, GFX1201); + MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX1250, GFX1250); MAP_COMPILER_STATE_TO_GFX12_TARGET(CK_TILE_ARCH_GFX12_GENERIC, GFX12_GENERIC); // Default to HOST @@ -708,6 +712,7 @@ CK_TILE_HOST auto hip_device_prop_gcn_arch_name_to_amdgcn_target(char const* tes MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX11_TARGET("gfx11_generic", GFX11_GENERIC); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX12_TARGET("gfx1200", GFX1200); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX12_TARGET("gfx1201", GFX1201); + MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX12_TARGET("gfx1250", GFX1250); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX12_TARGET("gfx12_generic", GFX12_GENERIC); // Default case diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index 06220d2780..8aa853c4bd 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -231,7 +231,7 @@ #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x00020000 #elif defined(__gfx101__) || defined(__gfx103__) // for GPU code #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31014000 -#elif defined(__gfx11__) || defined(__gfx12__) // for GPU code +#elif defined(__gfx11__) || defined(__gfx12__) || defined(__gfx1250__) // for GPU code #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31004000 #endif @@ -514,6 +514,12 @@ struct amdgcn_compiler_target_state static constexpr bool CK_TILE_ARCH_GFX1201 = false; #endif // __gfx1201__ +#if defined(__gfx1250__) + static constexpr bool CK_TILE_ARCH_GFX1250 = true; +#else + static constexpr bool CK_TILE_ARCH_GFX1250 = false; +#endif // __gfx1250__ + #if defined(__gfx12_generic__) static constexpr bool CK_TILE_ARCH_GFX12_GENERIC = true; #else @@ -570,6 +576,7 @@ CK_TILE_HOST_DEVICE static constexpr uint32_t count_values_of(T search, Ts... se amdgcn_compiler_target_state::CK_TILE_ARCH_GFX11_GENERIC, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1200, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1201, \ + amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1250, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX12_GENERIC // Sanity check: make sure only one target architecture is defined during device compile