From 1071198192162339e639bf55cc09ca109ae6a0c2 Mon Sep 17 00:00:00 2001 From: AviralGoelAMD Date: Wed, 23 Jul 2025 14:00:32 -0500 Subject: [PATCH] fix: add CK_GFX950_SUPPORT macro for gfx950 detection --- CMakeLists.txt | 2 ++ example/ck_tile/03_gemm/gemm_utils.hpp | 6 +++--- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6e032a30cf..da5a86523e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -236,6 +236,8 @@ endif() if (SUPPORTED_GPU_TARGETS MATCHES "gfx950") add_definitions(-DCK_USE_NATIVE_MX_SUPPORT) set(CK_USE_NATIVE_MX_SUPPORT "ON") + add_definitions(-DCK_GFX950_SUPPORT) + set(CK_GFX950_SUPPORT "ON") endif() option(CK_USE_FP8_ON_UNSUPPORTED_ARCH "Enable FP8 GEMM instances on older architectures" OFF) diff --git a/example/ck_tile/03_gemm/gemm_utils.hpp b/example/ck_tile/03_gemm/gemm_utils.hpp index 98db148e82..5068707160 100644 --- a/example/ck_tile/03_gemm/gemm_utils.hpp +++ b/example/ck_tile/03_gemm/gemm_utils.hpp @@ -21,7 +21,7 @@ template constexpr ck_tile::index_t get_k_warp_tile() { -#if defined(__gfx950__) +#if defined(CK_GFX950_SUPPORT) constexpr bool is_8bit_float = std::is_same_v || std::is_same_v; if constexpr(M_Warp_Tile == 32) @@ -38,7 +38,7 @@ constexpr ck_tile::index_t get_k_warp_tile() template constexpr ck_tile::index_t get_k_warp_tile_flatmm() { -#if defined(__gfx950__) +#if defined(CK_GFX950_SUPPORT) if constexpr(M_Warp_Tile == 32) return sizeof(PrecType) == 2 ? 16 : 64; else @@ -293,7 +293,7 @@ struct GemmConfigPreshuffle_3 : public GemmConfigBase static constexpr int kBlockPerCu = 2; static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Default; - static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_PRESHUFFLE_V3; + static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_PRESHUFFLE_V2; static constexpr bool Preshuffle = true; static constexpr bool DoubleSmemBuffer = true; };