ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-02 21:27:45 +00:00
Code Issues Packages Projects Releases Wiki Activity

debugging warp-level gemm classes

Browse Source
This commit is contained in:
Philip Maybank
2025-08-12 17:46:21 +01:00
parent 40ddc65a09
commit 55a1a331dc
2 changed files with 12 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -320,6 +320,17 @@ struct GfxConfig
return Gfx90x::kCM0PerLane;
}
}
static constexpr index_t get_k_cm1_per_lane()
{
if constexpr (GfxId == 1200)
{
return Gfx120x::kCM1PerLane;
}
else if constexpr (GfxId == 900)
{
return Gfx90x::kCM1PerLane;
}
}
};
} // namespace ck_tile

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

@@ -91,7 +91,7 @@ struct WarpGemmAttributeGenericImplF16F16F32M16N16K16
static constexpr index_t kCMLane = GfxConfig<GfxId>::get_k_cm_lane();
static constexpr index_t kCNLane = 16;
static constexpr index_t kCM0PerLane = GfxConfig<GfxId>::get_k_cm0_per_lane();
static constexpr index_t kCM1PerLane = 4;
static constexpr index_t kCM1PerLane = GfxConfig<GfxId>::get_k_cm1_per_lane();
// c_vec += a_vec * b_vec
template <bool post_nop_ = false>