diff --git a/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp b/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp index 1dcd62011a..23c4ad583e 100644 --- a/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp +++ b/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -73,7 +73,7 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_Base // for f16/bf16 sequence<2, 1>, // !! note here is different sequence<0, 0>>{}; - using WG = WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution; + using WG = WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution<>; constexpr auto c_block_dstr_encode = detail::make_embed_tile_distribution_encoding( c_block_outer_dstr_encoding, typename WG::CWarpDstrEncoding{}); diff --git a/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp b/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp index 0b812875c4..037bb7688c 100644 --- a/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp +++ b/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -49,7 +49,7 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_Base sequence<2, 1>, // !! note here is different sequence<0, 0>>{}; - using WG = WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution; + using WG = WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution<>; constexpr auto c_block_dstr_encode = detail::make_embed_tile_distribution_encoding( c_block_outer_dstr_encoding, typename WG::CWarpDstrEncoding{}); diff --git a/include/ck_tile/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_default_policy.hpp b/include/ck_tile/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_default_policy.hpp index 862fa0bbe3..cfbd78967f 100644 --- a/include/ck_tile/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_default_policy.hpp +++ b/include/ck_tile/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_default_policy.hpp @@ -47,11 +47,11 @@ struct BlockGemmASmemBSmemCRegV1DefaultPolicy if constexpr(NumWarp == 4 && kMPerBlock % 128 == 0 && kNPerBlock % 128 == 0 % kKPerBlock % 16 == 0) { - return make_tuple(WarpGemmMfmaF16F16F32M32N32K16{}, 2, 2); + return make_tuple(WarpGemmMfmaF16F16F32M32N32K16<>{}, 2, 2); } else { - return make_tuple(WarpGemmMfmaF16F16F32M32N32K16{}, 2, 2); + return make_tuple(WarpGemmMfmaF16F16F32M32N32K16<>{}, 2, 2); } #else using WG = WarpGemmMfmaDispatcher