From ad32373ff1d016dc045574c7b94edb0d122d596e Mon Sep 17 00:00:00 2001 From: joye Date: Wed, 21 May 2025 11:56:35 +0800 Subject: [PATCH] fix clang format --- example/ck_tile/03_gemm/universal_gemm.cpp | 3 ++- .../gemm_pipeline_ag_bg_cr_comp_async.hpp | 10 ++++---- ...emm_universal_pipeline_ag_bg_cr_policy.hpp | 24 +++++++++---------- 3 files changed, 19 insertions(+), 18 deletions(-) diff --git a/example/ck_tile/03_gemm/universal_gemm.cpp b/example/ck_tile/03_gemm/universal_gemm.cpp index b60a3b274b..9f1613f077 100644 --- a/example/ck_tile/03_gemm/universal_gemm.cpp +++ b/example/ck_tile/03_gemm/universal_gemm.cpp @@ -199,7 +199,8 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& ck_tile::integral_constant{}, ck_tile::integral_constant{}); -#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V4) +#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V4 || \ + CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_ASYNC) if(tail_num == ck_tile::TailNumber::Three) { RunSplitk( diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp index 150e6fe23a..e14591688b 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp @@ -331,11 +331,11 @@ struct GemmPipelineAgBgCrCompAsync : public BaseGemmPipelineAgBgCrCompAsync) && - !(is_tile_window_linear_v) && - !(is_tile_window_linear_v) && - !(is_tile_window_linear_v), - "LDS windows must not be linear"); + static_assert( + !(is_tile_window_linear_v)&&!(is_tile_window_linear_v)&&!( + is_tile_window_linear_v< + decltype(b_lds_ld_window0)>)&&!(is_tile_window_linear_v), + "LDS windows must not be linear"); buffer_load_fence(a_number_of_access + b_number_of_access); block_sync_lds(); diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp index b31e5dfe48..7aab2ec27a 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp @@ -177,10 +177,10 @@ struct UniversalGemmBasePolicy constexpr index_t BlockSize = Problem::kBlockSize; constexpr index_t VecLoadSize = GetVectorSizeB(); using TileEncodingPattern = TileDistributionEncodingPattern2D; + KPerBlock, + NPerBlock, + VecLoadSize, + BTileAccessPattern>; constexpr auto BK0 = number{}; constexpr auto BK1 = number{}; @@ -621,14 +621,14 @@ struct UniversalGemmPipelineAgBgCrPolicy using BlockWarps = typename Problem::BlockGemmShape::BlockWarps; using WarpTile = typename Problem::BlockGemmShape::WarpTile; using WarpGemm = WarpGemmMfmaDispatcher; + typename Problem::ComputeDataType, + typename Problem::CDataType, + WarpTile::at(I0), + WarpTile::at(I1), + WarpTile::at(I2), + Problem::TransposeC, + false, + Problem::UseStructuredSparsity>; using BlockGemmPolicy = BlockGemmASmemBSmemCRegV1CustomPolicy