From 5cd27923503e165872ed7c2113bc73f330a4b2d5 Mon Sep 17 00:00:00 2001 From: Philip Maybank Date: Fri, 15 Aug 2025 10:07:45 +0100 Subject: [PATCH] debug build-time error by setting kABK etc - still getting run-time error --- .../00_add_vector/add_vector.cpp | 12 +- include/ck_tile/core/arch/arch.hpp | 114 +++++++++++++----- include/ck_tile/ops/gemm/warp/warp_gemm.hpp | 4 +- ...ribute_generic_impl_F16F16F32M16N16K16.hpp | 12 +- ...attribute_wmma_impl_F16F16F32M16N16K16.hpp | 76 ++++++++++++ 5 files changed, 174 insertions(+), 44 deletions(-) create mode 100644 include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_F16F16F32M16N16K16.hpp diff --git a/example/ck_tile/99_toy_example/00_add_vector/add_vector.cpp b/example/ck_tile/99_toy_example/00_add_vector/add_vector.cpp index 4431aadf71..ae82d50138 100644 --- a/example/ck_tile/99_toy_example/00_add_vector/add_vector.cpp +++ b/example/ck_tile/99_toy_example/00_add_vector/add_vector.cpp @@ -75,20 +75,20 @@ bool run(const ck_tile::ArgParser& arg_parser) // These will hold the *values* for the ck_tile::sequence types // They are initialized based on the GfxId - constexpr ck_tile::index_t selected_warp_tile = (GfxId == 1200) ? Gfx120x::WarpTile - : (GfxId == 900) ? Gfx90x::WarpTile + constexpr ck_tile::index_t selected_warp_tile = (GfxId == 1201) ? Gfx1201::WarpTile + : (GfxId == 900) ? Gfx900::WarpTile : /* else */ Generic::WarpTile; // Use if constexpr to select the compile-time constants for the current GfxId bool fail = false; - if constexpr(GfxId == 1200) + if constexpr(GfxId == 1201) { - std::cout << "Using gfx120x-optimized parameters (template specialization)." << std::endl; + std::cout << "Using gfx1201-optimized parameters (template specialization)." << std::endl; } else if constexpr(GfxId == 900) { - std::cout << "Using gfx90x-optimized parameters (template specialization)." << std::endl; + std::cout << "Using gfx900-optimized parameters (template specialization)." << std::endl; } else { // Fallback for GfxId == 0 or unknown @@ -201,7 +201,7 @@ int main(int argc, char* argv[]) std::string arch_name = props.gcnArchName; if(data_type == "fp16" && (arch_name.find("gfx12") != std::string::npos)) - return run(arg_parser) ? 0 : -2; + return run(arg_parser) ? 0 : -2; else if(data_type == "fp16" && (arch_name.find("gfx908") != std::string::npos)) return run(arg_parser) ? 0 : -2; else diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp index 86de1b8787..ad494c1b79 100644 --- a/include/ck_tile/core/arch/arch.hpp +++ b/include/ck_tile/core/arch/arch.hpp @@ -239,32 +239,42 @@ CK_TILE_HOST_DEVICE constexpr index_t get_smem_capacity() // We'll define all parameters for all supported architectures here. // ============================================================================ -// Parameters for gfx90x (example values, adjust as needed) -namespace Gfx90x { +// Parameters for gfx900 (example values, adjust as needed) +namespace Gfx900 { constexpr ck_tile::index_t WarpTile = 64; constexpr ck_tile::index_t VecLenFP16 = 4; constexpr ck_tile::index_t VecLenFP32 = 4; +constexpr ck_tile::index_t kAMLane = 16; +constexpr ck_tile::index_t kBNLane = 16; +constexpr ck_tile::index_t kABKLane = 4; +constexpr ck_tile::index_t kABKPerLane = 4; + constexpr ck_tile::index_t kCMLane = 4; constexpr ck_tile::index_t kCM0PerLane = 1; constexpr ck_tile::index_t kCM1PerLane = 4; } -// Parameters for gfx120x (using a namespace for organization or just global constexpr) -namespace Gfx120x { +// Parameters for gfx1201 (using a namespace for organization or just global constexpr) +namespace Gfx1201 { constexpr ck_tile::index_t WarpTile = 32; -constexpr ck_tile::index_t VecLenFP16 = 16; +constexpr ck_tile::index_t VecLenFP16 = 8; constexpr ck_tile::index_t VecLenFP32 = 8; +constexpr ck_tile::index_t kAMLane = 16; +constexpr ck_tile::index_t kBNLane = 16; +constexpr ck_tile::index_t kABKLane = 2; +constexpr ck_tile::index_t kABKPerLane = 8; + constexpr ck_tile::index_t kCMLane = 2; -constexpr ck_tile::index_t kCM0PerLane = 8; -constexpr ck_tile::index_t kCM1PerLane = 1; +constexpr ck_tile::index_t kCM0PerLane = 1; +constexpr ck_tile::index_t kCM1PerLane = 8; } // Generic Parameters - should never be used in this example -// templated run function should only be instantiated for Gfx120x and Gfx90x +// templated run function should only be instantiated for Gfx1201 and Gfx900 namespace Generic { constexpr ck_tile::index_t WarpTile = -1; } @@ -275,62 +285,106 @@ struct GfxConfig { static constexpr index_t get_vec_len_fp16() { - if constexpr (GfxId == 1200) + if constexpr (GfxId == 1201) { - return Gfx120x::VecLenFP16; + return Gfx1201::VecLenFP16; } else if constexpr (GfxId == 900) { - return Gfx90x::VecLenFP16; + return Gfx900::VecLenFP16; } } static constexpr index_t get_vec_len_fp32() { - if constexpr (GfxId == 1200) + if constexpr (GfxId == 1201) { - return Gfx120x::VecLenFP32; + return Gfx1201::VecLenFP32; } else if constexpr (GfxId == 900) { - return Gfx90x::VecLenFP32; + return Gfx900::VecLenFP32; } - } + } static constexpr index_t get_k_cm_lane() { - if constexpr (GfxId == 1200) + if constexpr (GfxId == 1201) { - return Gfx120x::kCMLane; + return Gfx1201::kCMLane; } else if constexpr (GfxId == 900) { - return Gfx90x::kCMLane; - } + return Gfx900::kCMLane; + } } - + static constexpr index_t get_k_cm0_per_lane() { - if constexpr (GfxId == 1200) + if constexpr (GfxId == 1201) { - return Gfx120x::kCM0PerLane; + return Gfx1201::kCM0PerLane; } else if constexpr (GfxId == 900) { - return Gfx90x::kCM0PerLane; - } - } + return Gfx900::kCM0PerLane; + } + } static constexpr index_t get_k_cm1_per_lane() { - if constexpr (GfxId == 1200) + if constexpr (GfxId == 1201) { - return Gfx120x::kCM1PerLane; + return Gfx1201::kCM1PerLane; } else if constexpr (GfxId == 900) { - return Gfx90x::kCM1PerLane; - } - } + return Gfx900::kCM1PerLane; + } + } + static constexpr index_t get_k_am_lane() + { + if constexpr (GfxId == 1201) + { + return Gfx1201::kAMLane; + } + else if constexpr (GfxId == 900) + { + return Gfx900::kAMLane; + } + } + static constexpr index_t get_k_bn_lane() + { + if constexpr (GfxId == 1201) + { + return Gfx1201::kBNLane; + } + else if constexpr (GfxId == 900) + { + return Gfx900::kBNLane; + } + } + static constexpr index_t get_k_abk_lane() + { + if constexpr (GfxId == 1201) + { + return Gfx1201::kABKLane; + } + else if constexpr (GfxId == 900) + { + return Gfx900::kABKLane; + } + } + static constexpr index_t get_k_abk_per_lane() + { + if constexpr (GfxId == 1201) + { + return Gfx1201::kABKPerLane; + } + else if constexpr (GfxId == 900) + { + return Gfx900::kABKPerLane; + } + } }; } // namespace ck_tile diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm.hpp index 7eb665f59c..c95a5d86fb 100644 --- a/include/ck_tile/ops/gemm/warp/warp_gemm.hpp +++ b/include/ck_tile/ops/gemm/warp/warp_gemm.hpp @@ -35,7 +35,7 @@ using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl< WarpGemmAttributeGeneric>>; using WarpGemmWmmaF16F16F32M16N16K16 = WarpGemmImpl< - WarpGemmAttributeGeneric>>; + WarpGemmAttributeGeneric>>; #if 0 #if defined(__gfx950__) @@ -92,7 +92,7 @@ using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution = using WarpGemmWmmaF16F16F32M16N16K16TransposedCDistribution = WarpGemmImpl>>; + WarpGemmAttributeGenericImplF16F16F32M16N16K16<1201, WGAttrCtlEnum::Default_>>>; #if 0 diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp index 6a52b9d3eb..547e0f9670 100644 --- a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp +++ b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp @@ -71,7 +71,7 @@ struct WarpGemmAttributeGenericImplF16F16F32M16N16K16 using CDataType = float; static constexpr index_t VecLenFP16 = GfxConfig::get_vec_len_fp16(); - static constexpr index_t VecLenFP32 = GfxConfig::get_vec_len_fp32(); + static constexpr index_t VecLenFP32 = GfxConfig::get_vec_len_fp32(); using AVecType = ext_vector_t; using BVecType = ext_vector_t; using CVecType = ext_vector_t; @@ -85,13 +85,13 @@ struct WarpGemmAttributeGenericImplF16F16F32M16N16K16 static constexpr index_t kAMLane = 16; static constexpr index_t kBNLane = 16; - static constexpr index_t kABKLane = 4; - static constexpr index_t kABKPerLane = 4; + static constexpr index_t kABKLane = GfxConfig::get_k_abk_lane(); // 4 for gfx9, 2 for gfx12 + static constexpr index_t kABKPerLane = GfxConfig::get_k_abk_per_lane(); // 4 for gfx9, 8 for gfx12; - static constexpr index_t kCMLane = GfxConfig::get_k_cm_lane(); + static constexpr index_t kCMLane = GfxConfig::get_k_cm_lane(); // 4 for gfx9, 2 for gfx12 static constexpr index_t kCNLane = 16; - static constexpr index_t kCM0PerLane = GfxConfig::get_k_cm0_per_lane(); - static constexpr index_t kCM1PerLane = GfxConfig::get_k_cm1_per_lane(); + static constexpr index_t kCM0PerLane = 1; + static constexpr index_t kCM1PerLane = GfxConfig::get_k_cm1_per_lane(); // 4 for gfx9, 8 for gfx12 // c_vec += a_vec * b_vec template diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_F16F16F32M16N16K16.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_F16F16F32M16N16K16.hpp new file mode 100644 index 0000000000..22a11265d3 --- /dev/null +++ b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_F16F16F32M16N16K16.hpp @@ -0,0 +1,76 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include "ck_tile/core.hpp" +#include "ck_tile/core/config.hpp" + +namespace ck_tile { + +struct WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16_gfx12 +{ + using ADataType = fp16_t; + using BDataType = fp16_t; + using CDataType = float; + + using AVecType = ext_vector_t; + using BVecType = ext_vector_t; + using CVecType = ext_vector_t; + + static constexpr index_t kM = 16; + static constexpr index_t kN = 16; + static constexpr index_t kK = 16; + + static constexpr index_t kAMLane = 16; + static constexpr index_t kBNLane = 16; + static constexpr index_t kABK0PerLane = 1; + static constexpr index_t kABKLane = 2; + static constexpr index_t kABK1PerLane = 8; + + static constexpr index_t kCMLane = 2; + static constexpr index_t kCNLane = 16; + static constexpr index_t kCM0PerLane = 1; + static constexpr index_t kCM1PerLane = 8; + + using kABPs2RHssMajor = sequence<2, 1>; + using kABPs2RHssMinor = sequence<1, 0>; + using kABYs2RHsMajor = sequence<2, 2>; + using kABYs2RHsMinor = sequence<0, 2>; + + using kCPs2RHssMajor = sequence<2, 1>; + using kCPs2RHssMinor = sequence<1, 0>; + using kCYs2RHsMajor = sequence<2, 2>; + using kCYs2RHsMinor = sequence<0, 2>; + // c_vec += a_vec * b_vec + template + CK_TILE_DEVICE void operator()(CVecType& c_vec, + const AVecType& a_vec, + const BVecType& b_vec, + bool_constant = {}) const + { +#if defined(__gfx12__) + c_vec = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_vec, b_vec, c_vec); +#else + ck_tile::ignore = c_vec; + ck_tile::ignore = a_vec; + ck_tile::ignore = b_vec; +#endif + } + + // c_vec = a_vec * b_vec + template + CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const + { +#if defined(__gfx12__) + return bit_cast( + __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_vec, b_vec, fp32x8_t{0.f})); +#else + ck_tile::ignore = a_vec; + ck_tile::ignore = b_vec; + return CVecType{0.f}; +#endif + } +}; + +} \ No newline at end of file