ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 17:51:40 +00:00
Code Issues Packages Projects Releases Wiki Activity

debug build-time error by setting kABK etc - still getting run-time error

Browse Source
This commit is contained in:
Philip Maybank
2025-08-15 10:07:45 +01:00
parent 55a1a331dc
commit 5cd2792350
5 changed files with 174 additions and 44 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
example/ck_tile/99_toy_example/00_add_vector/add_vector.cpp
View File

@@ -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<ck_tile::half_t, 1200>(arg_parser) ? 0 : -2;
return run<ck_tile::half_t, 1201>(arg_parser) ? 0 : -2;
else if(data_type == "fp16" && (arch_name.find("gfx908") != std::string::npos))
return run<ck_tile::half_t, 900>(arg_parser) ? 0 : -2;
else

114
include/ck_tile/core/arch/arch.hpp
View File

@@ -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

4
include/ck_tile/ops/gemm/warp/warp_gemm.hpp
View File

@@ -35,7 +35,7 @@ using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl<
WarpGemmAttributeGeneric<WarpGemmAttributeGenericImplF16F16F32M16N16K16<900, WGAttrCtlEnum::Default_>>>;
using WarpGemmWmmaF16F16F32M16N16K16 = WarpGemmImpl<
WarpGemmAttributeGeneric<WarpGemmAttributeGenericImplF16F16F32M16N16K16<1200, WGAttrCtlEnum::Default_>>>;
WarpGemmAttributeGeneric<WarpGemmAttributeGenericImplF16F16F32M16N16K16<1201, WGAttrCtlEnum::Default_>>>;
#if 0
#if defined(__gfx950__)
@@ -92,7 +92,7 @@ using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution =
using WarpGemmWmmaF16F16F32M16N16K16TransposedCDistribution =
WarpGemmImpl<WarpGemmAttributeGenericTransposedCDistribution<
WarpGemmAttributeGenericImplF16F16F32M16N16K16<1200, WGAttrCtlEnum::Default_>>>;
WarpGemmAttributeGenericImplF16F16F32M16N16K16<1201, WGAttrCtlEnum::Default_>>>;
#if 0

12
include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp
View File

@@ -71,7 +71,7 @@ struct WarpGemmAttributeGenericImplF16F16F32M16N16K16
using CDataType = float;
static constexpr index_t VecLenFP16 = GfxConfig<GfxId>::get_vec_len_fp16();
static constexpr index_t VecLenFP32 = GfxConfig<GfxId>::get_vec_len_fp32();
static constexpr index_t VecLenFP32 = GfxConfig<GfxId>::get_vec_len_fp32();
using AVecType = ext_vector_t<fp16_t, VecLenFP16>;
using BVecType = ext_vector_t<fp16_t, VecLenFP16>;
using CVecType = ext_vector_t<float, VecLenFP32>;
@@ -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<GfxId>::get_k_abk_lane(); // 4 for gfx9, 2 for gfx12
static constexpr index_t kABKPerLane = GfxConfig<GfxId>::get_k_abk_per_lane(); // 4 for gfx9, 8 for gfx12;
static constexpr index_t kCMLane = GfxConfig<GfxId>::get_k_cm_lane();
static constexpr index_t kCMLane = GfxConfig<GfxId>::get_k_cm_lane(); // 4 for gfx9, 2 for gfx12
static constexpr index_t kCNLane = 16;
static constexpr index_t kCM0PerLane = GfxConfig<GfxId>::get_k_cm0_per_lane();
static constexpr index_t kCM1PerLane = GfxConfig<GfxId>::get_k_cm1_per_lane();
static constexpr index_t kCM0PerLane = 1;
static constexpr index_t kCM1PerLane = GfxConfig<GfxId>::get_k_cm1_per_lane(); // 4 for gfx9, 8 for gfx12
// c_vec += a_vec * b_vec
template <bool post_nop_ = false>

76
include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_F16F16F32M16N16K16.hpp Normal file
View File

@@ -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<fp16_t, 8>;
using BVecType = ext_vector_t<fp16_t, 8>;
using CVecType = ext_vector_t<float, 8>;
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 <bool clamp = false, bool post_nop_ = false>
CK_TILE_DEVICE void operator()(CVecType& c_vec,
const AVecType& a_vec,
const BVecType& b_vec,
bool_constant<post_nop_> = {}) 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 <bool clamp = false>
CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const
{
#if defined(__gfx12__)
return bit_cast<CVecType>(
__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
}
};
}