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

Disable SMFMA gfx90a (#2184)

Browse Source 
* sparsity fix for gfx90a

* reverting tile_engine changes
This commit is contained in:
Khushbu Agarwal
2025-05-12 09:56:23 -07:00
committed by GitHub
parent b49f7de81f
commit f05e45ba59
3 changed files with 3 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -109,20 +109,11 @@ using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK
4>>;
// fp16 2:4 structured sparsity
#if defined(__gfx94__) || defined(__gfx95__)
using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<
WarpGemmAttributeSmfmacImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;
using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<
WarpGemmAttributeSmfmacImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;
#else // gfx 90a does not support smfmac
using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,
2>>;
using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,
2>>;
#endif
// bf16
using WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl<

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

@@ -49,7 +49,7 @@ struct WarpGemmAttributeSmfmacImplF16F16F32M32N32K16
const int32_t& idx,
bool_constant<post_nop_> = {}) const
{
#if defined(__gfx9__)
#if defined(__gfx94_) or defined(__gfx95_)
c_vec = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a_vec, b_vec, c_vec, idx, 0, 0);
#else
ck_tile::ignore = c_vec;
@@ -100,7 +100,7 @@ struct WarpGemmAttributeSmfmacImplF16F16F32M16N16K32
const int32_t& idx,
bool_constant<post_nop_> = {}) const
{
#if defined(__gfx9__)
#if defined(__gfx94_) or defined(__gfx95_)
c_vec = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a_vec, b_vec, c_vec, idx, 0, 0);
#else
ck_tile::ignore = c_vec;