diff --git a/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp b/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp index 5f6c56d0e1..c6fbfb4452 100644 --- a/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp +++ b/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp @@ -26,7 +26,8 @@ struct BlockGemmARegBSmemCRegV1K8Policy #if !defined(TOY_FA_FWD_QK_SWIZZLE) return make_tuple(WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution{}, 4, 1); #else - return make_tuple(WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution{}, 4, 1); + return make_tuple( + WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution{}, 4, 1); #endif } else diff --git a/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp b/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp index 32acb0f1b6..cfbd7d6376 100644 --- a/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp +++ b/example/ck_tile/99_toy_example/03_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp @@ -374,9 +374,9 @@ struct BlockGemmPipelineAGmemBGmemCReg { return operator()( a_dram_block_window_tmp, - [](const ADataType& a) { return a; }, + [](const ADataType & a) { return a; }, b_dram_block_window_tmp, - [](const BDataType& b) { return b; }, + [](const BDataType & b) { return b; }, a_reg_block_tensor_tmp, p_smem); } @@ -388,7 +388,7 @@ struct BlockGemmPipelineAGmemBGmemCReg { return operator()( b_dram_block_window_tmp, - [](const BDataType& b) { return b; }, + [](const BDataType & b) { return b; }, a_reg_block_tensor_tmp, p_smem); } diff --git a/example/ck_tile/99_toy_example/03_flash_attention_fwd/flash_attention_fwd_impl.hpp b/example/ck_tile/99_toy_example/03_flash_attention_fwd/flash_attention_fwd_impl.hpp index 90779420c1..fced47a536 100644 --- a/example/ck_tile/99_toy_example/03_flash_attention_fwd/flash_attention_fwd_impl.hpp +++ b/example/ck_tile/99_toy_example/03_flash_attention_fwd/flash_attention_fwd_impl.hpp @@ -64,9 +64,9 @@ struct FlashAttentionFwdImpl constexpr index_t kNPerBlock = kN1PerBlock; constexpr index_t kKPerBlock = kK1PerBlock; #if !defined(TOY_FA_FWD_QK_SWIZZLE) - constexpr index_t kKPack = 4; + constexpr index_t kKPack = 4; #else - constexpr index_t kKPack = 8; + constexpr index_t kKPack = 8; #endif constexpr auto dataTypeSize = sizeof(VDataType); diff --git a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp index 5f6c56d0e1..c6fbfb4452 100644 --- a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp +++ b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp @@ -26,7 +26,8 @@ struct BlockGemmARegBSmemCRegV1K8Policy #if !defined(TOY_FA_FWD_QK_SWIZZLE) return make_tuple(WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution{}, 4, 1); #else - return make_tuple(WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution{}, 4, 1); + return make_tuple( + WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution{}, 4, 1); #endif } else diff --git a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp index 32acb0f1b6..cfbd7d6376 100644 --- a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp +++ b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/block_gemm_pipeline_agmem_bgmem_creg_v2_askiplds.hpp @@ -374,9 +374,9 @@ struct BlockGemmPipelineAGmemBGmemCReg { return operator()( a_dram_block_window_tmp, - [](const ADataType& a) { return a; }, + [](const ADataType & a) { return a; }, b_dram_block_window_tmp, - [](const BDataType& b) { return b; }, + [](const BDataType & b) { return b; }, a_reg_block_tensor_tmp, p_smem); } @@ -388,7 +388,7 @@ struct BlockGemmPipelineAGmemBGmemCReg { return operator()( b_dram_block_window_tmp, - [](const BDataType& b) { return b; }, + [](const BDataType & b) { return b; }, a_reg_block_tensor_tmp, p_smem); } diff --git a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp index 073b4eb6aa..26e8d02845 100644 --- a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp +++ b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp @@ -90,13 +90,13 @@ struct FlashAttentionFwd const auto f = [](index_t dividend, index_t divisor) { index_t quotient = dividend / divisor; index_t modulus = dividend - quotient * divisor; - + return make_tuple(quotient, modulus); }; - + const auto [itmp, id_tile_n] = f(id_block, num_tile_n1); const auto [id_tile_batch, id_tile_m] = f(itmp, num_tile_m0); - + const index_t iBatch = __builtin_amdgcn_readfirstlane(id_tile_batch); const index_t iM0 = __builtin_amdgcn_readfirstlane(id_tile_m * kM0PerBlock); const index_t iN1 = __builtin_amdgcn_readfirstlane(id_tile_n * kN1PerBlock); diff --git a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp index 90779420c1..fced47a536 100644 --- a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp +++ b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp @@ -64,9 +64,9 @@ struct FlashAttentionFwdImpl constexpr index_t kNPerBlock = kN1PerBlock; constexpr index_t kKPerBlock = kK1PerBlock; #if !defined(TOY_FA_FWD_QK_SWIZZLE) - constexpr index_t kKPack = 4; + constexpr index_t kKPack = 4; #else - constexpr index_t kKPack = 8; + constexpr index_t kKPack = 8; #endif constexpr auto dataTypeSize = sizeof(VDataType); diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp index 3489d6f9a1..38e93a5425 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy.hpp @@ -948,7 +948,7 @@ struct BlockFmhaPipelineQXKSVSCustomPolicy : BlockFmhaPipelineQXCustomPolicy{}; // return - // WarpGemmImpl>>{}; } diff --git a/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp b/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp index 0c8baaf191..9b234b5797 100644 --- a/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp +++ b/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.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 @@ -568,7 +568,7 @@ struct FusedMoeGemmPipelineFlatmmPolicy std::is_same_v && S_::Warp_M0 == 32 && S_::Warp_N0 == 32 && S_::Warp_K0 == 16) { - return WarpGemmImpl, 2>>{}; } @@ -576,7 +576,7 @@ struct FusedMoeGemmPipelineFlatmmPolicy std::is_same_v && S_::Warp_M0 == 32 && S_::Warp_N0 == 32 && S_::Warp_K0 == 32) { - return WarpGemmImpl, 2>>{}; } @@ -695,7 +695,7 @@ struct FusedMoeGemmPipelineFlatmmPolicy std::is_same_v && S_::Warp_M0 == 32 && S_::Warp_N0 == 32 && S_::Warp_K0 == 16) { - return WarpGemmImpl, 2>>{}; } @@ -703,7 +703,7 @@ struct FusedMoeGemmPipelineFlatmmPolicy std::is_same_v && S_::Warp_M0 == 32 && S_::Warp_N0 == 32 && S_::Warp_K0 == 32) { - return WarpGemmImpl, 2>>{}; } diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm.hpp index 47b91ccbf7..4a7f0e344b 100644 --- a/include/ck_tile/ops/gemm/warp/warp_gemm.hpp +++ b/include/ck_tile/ops/gemm/warp/warp_gemm.hpp @@ -15,19 +15,19 @@ namespace ck_tile { // fp16 using WarpGemmMfmaF16F16F32M32N32K8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; #if defined(__gfx950__) template using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; #else template -using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl, 2, AttrNumAccess>>; @@ -36,42 +36,42 @@ using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; #else template -using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl, 2, AttrNumAccess>>; #endif -using WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl, 1>>; -using WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl, 2>>; using WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution = - WarpGemmImpl>>; using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution = - WarpGemmImpl>>; #if defined(__gfx950__) template using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution = - WarpGemmImpl, AttrNumAccess>>; #else template using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution = - WarpGemmImpl, 2, AttrNumAccess>>; @@ -80,13 +80,13 @@ using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution = #if defined(__gfx950__) template using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution = - WarpGemmImpl, AttrNumAccess>>; #else template using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution = - WarpGemmImpl, 2, AttrNumAccess>>; @@ -94,32 +94,32 @@ using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution = #if defined(__gfx950__) using WarpGemmMfmaF16F16F32M16N16K32SwizzleBTransposedCDistribution = - WarpGemmImpl, 1>>; using WarpGemmMfmaBf16Bf16F32M16N16K32SwizzleBTransposedCDistribution = - WarpGemmImpl, 1>>; #endif #if defined(__gfx950__) using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution = - WarpGemmImpl>>; #else using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution = - WarpGemmImpl, 2>>; #endif -using WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl, 4>>; -using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl, 4>>; @@ -132,19 +132,19 @@ using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl>>; + WarpGemmAttributeMfma>>; using WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; #if defined(__gfx950__) template using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; #else template -using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl, 2, AttrNumAccess>>; @@ -153,43 +153,43 @@ using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; #else template -using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl, 2, AttrNumAccess>>; #endif -using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl, 1>>; using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA = - WarpGemmImpl, 2>>; using WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution = - WarpGemmImpl>>; using WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution = - WarpGemmImpl>>; #if defined(__gfx950__) template using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution = - WarpGemmImpl, AttrNumAccess>>; #else template using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution = - WarpGemmImpl, 2, AttrNumAccess>>; @@ -198,13 +198,13 @@ using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution = #if defined(__gfx950__) template using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution = - WarpGemmImpl, AttrNumAccess>>; #else template using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution = - WarpGemmImpl, 2, AttrNumAccess>>; @@ -212,127 +212,127 @@ using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution = #if defined(__gfx950__) using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution = - WarpGemmImpl>>; #else using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution = - WarpGemmImpl, 2>>; #endif -using WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl, 4>>; -using WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl, 4>>; // fp8 using WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; -using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl, 2>>; -using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl, 2>>; using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; -using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl, 2>>; -using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl, 2>>; -using WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl>>; -using WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl>>; -using WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl>>; -using WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl>>; template using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; template using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; template using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; template using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl< - WarpGemmAtrributeMfma, + WarpGemmAttributeMfma, AttrNumAccess>>; using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed = - WarpGemmImpl>>; using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed = - WarpGemmImpl>>; using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed = - WarpGemmImpl>>; using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed = - WarpGemmImpl>>; template using WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution = - WarpGemmImpl, 2, swizzle_factor>>; // int8 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed = - WarpGemmImpl>>; using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl< - WarpGemmAtrributeMfma>>; + WarpGemmAttributeMfma>>; using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed = - WarpGemmImpl>>; } // namespace ck_tile diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp index 97fab489ab..36a9955912 100644 --- a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp +++ b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp @@ -19,7 +19,7 @@ enum class WGAttrNumAccessEnum template -struct WarpGemmAtrributeMfma +struct WarpGemmAttributeMfma { using Impl = remove_cvref_t; static constexpr auto AttrNumAccess = AttrNumAccess_; @@ -103,7 +103,7 @@ struct WarpGemmAtrributeMfma template -struct WarpGemmAtrributeMfmaIterateK +struct WarpGemmAttributeMfmaIterateK { static_assert(kKIter > 0, "wrong!"); @@ -367,7 +367,7 @@ struct WarpGemmAtrributeMfmaIterateK template -struct WarpGemmAtrributeMfmaTransposedCDistribution +struct WarpGemmAttributeMfmaTransposedCDistribution { using Impl = remove_cvref_t; static constexpr auto AttrNumAccess = AttrNumAccess_; @@ -450,7 +450,7 @@ struct WarpGemmAtrributeMfmaTransposedCDistribution }; template -struct WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB +struct WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB { using Impl = remove_cvref_t; @@ -546,7 +546,7 @@ struct WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB template -struct WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution +struct WarpGemmAttributeMfmaIterateKAndTransposedCDistribution { using Impl = remove_cvref_t; static constexpr auto AttrNumAccess = AttrNumAccess_; @@ -574,13 +574,13 @@ struct WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution CK_TILE_DEVICE static constexpr auto get_awarp_dstr_encoding() { - return WarpGemmAtrributeMfmaIterateK:: + return WarpGemmAttributeMfmaIterateK:: get_bwarp_dstr_encoding(); } CK_TILE_DEVICE static constexpr auto get_bwarp_dstr_encoding() { - return WarpGemmAtrributeMfmaIterateK:: + return WarpGemmAttributeMfmaIterateK:: get_awarp_dstr_encoding(); } @@ -696,7 +696,7 @@ struct WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution }; template -struct WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB +struct WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB { using Impl = remove_cvref_t; @@ -840,7 +840,7 @@ struct WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB }; template -struct WarpGemmAtrributeMfmaIterateK_SwizzleA +struct WarpGemmAttributeMfmaIterateK_SwizzleA { using Impl = remove_cvref_t;