ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-30 11:47:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add get_device_ip()

Browse Source 
Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>
This commit is contained in:
TianyuanWu
2025-08-06 10:30:00 +08:00
parent eb780fb003
commit 2811a4e1de
2 changed files with 13 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -246,4 +246,13 @@ struct gfx12_t
{
};
CK_TILE_DEVICE static constexpr auto get_device_ip()
{
#if defined(__gfx11__)
return gfx11_t{};
#else // if defined(__gfx12__)
return gfx12_t{};
#endif
}
} // namespace ck_tile

19
include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp
View File

@@ -76,12 +76,9 @@ struct WarpGemmAttributeWmmaImpl
}
};
using DeviceIp = remove_cvref_t<decltype(ck_tile::get_device_ip())>;
using WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16 = WarpGemmAttributeWmmaImpl<WmmaTraits<
#if defined(__gfx11__)
gfx11_t,
#else // if defined(__gfx12__)
gfx12_t,
#endif
DeviceIp,
fp16_t,
fp16_t,
float,
@@ -90,11 +87,7 @@ using WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16 = WarpGemmAttributeWmmaImpl
16>>;
using WarpGemmAttributeWmmaImpl_f32_16x16x16_bf16_bf16 = WarpGemmAttributeWmmaImpl<WmmaTraits<
#if defined(__gfx11__)
gfx11_t,
#else // if defined(__gfx12__)
gfx12_t,
#endif
DeviceIp,
bf16_t,
bf16_t,
float,
@@ -103,11 +96,7 @@ using WarpGemmAttributeWmmaImpl_f32_16x16x16_bf16_bf16 = WarpGemmAttributeWmmaIm
16>>;
using WarpGemmAttributeWmmaImpl_i32_16x16x16_i8_i8 = WarpGemmAttributeWmmaImpl<WmmaTraits<
#if defined(__gfx11__)
gfx11_t,
#else // if defined(__gfx12__)
gfx12_t,
#endif
DeviceIp,
int8_t,
int8_t,
int32_t,