ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 02:02:46 +00:00
Code Issues Packages Projects Releases Wiki Activity

fix:tf32:fix build fail for all supported targets (#2942)

Browse Source 
* fix:tf32:fix build fail for all supported targets

* new fix code

[ROCm/composable_kernel commit: 0f04f020d9]
This commit is contained in:
yinglu
2025-09-29 23:04:11 +08:00
committed by GitHub
parent b6cb76a555
commit f9daaa9724
2 changed files with 19 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
CMakeLists.txt
View File

@@ -220,6 +220,9 @@ rocm_check_target_ids(SUPPORTED_GPU_TARGETS
message(STATUS "Building CK for the following targets: ${SUPPORTED_GPU_TARGETS}")
# Cache SUPPORTED_GPU_TARGETS for debug
set(SUPPORTED_GPU_TARGETS "${SUPPORTED_GPU_TARGETS}" CACHE STRING "List of supported GPU targets")
if (SUPPORTED_GPU_TARGETS MATCHES "gfx9|gfx11|gfx12")
message(STATUS "Enabling XDL instances")
add_definitions(-DCK_USE_XDL)

16
include/ck/tensor_operation/gpu/warp/xdlops_gemm.hpp
View File

@@ -1277,13 +1277,29 @@ struct MfmaSelector
template <>
constexpr auto GetMfma<tf32_t, 32, 32>()
{
#if defined(__gfx12__)
return MfmaInstr::wmma_unsupport_16x16_gfx12;
#elif defined(__gfx11__)
return MfmaInstr::wmma_unsupport_16x16_gfx11;
#elif defined(__gfx942__)
return MfmaInstr::mfma_f32_32x32x4xf32;
#else
return MfmaInstr::mfma_f32_32x32x2f32;
#endif
}
template <>
constexpr auto GetMfma<tf32_t, 16, 16>()
{
#if defined(__gfx12__)
return MfmaInstr::wmma_unsupport_16x16_gfx12;
#elif defined(__gfx11__)
return MfmaInstr::wmma_unsupport_16x16_gfx11;
#elif defined(__gfx942__)
return MfmaInstr::mfma_f32_16x16x8xf32;
#else
return MfmaInstr::mfma_f32_16x16x4f32;
#endif
}
template <>