From ebda5006f84f02672889089444decf79b62f8ed7 Mon Sep 17 00:00:00 2001 From: Clement Lin Date: Sun, 18 May 2025 18:34:17 +0800 Subject: [PATCH] Fix clang format --- .../block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp | 3 ++- .../03_flash_attention_fwd/flash_attention_fwd_impl.hpp | 6 +++--- .../block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp | 3 ++- .../04_codegen_flash_attention_fwd/flash_attention_fwd.hpp | 6 +++--- .../flash_attention_fwd_impl.hpp | 6 +++--- 5 files changed, 13 insertions(+), 11 deletions(-) diff --git a/example/ck_tile/tutorial/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp b/example/ck_tile/tutorial/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp index d7e6ee4c96..9943fdbef4 100644 --- a/example/ck_tile/tutorial/03_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp +++ b/example/ck_tile/tutorial/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/tutorial/03_flash_attention_fwd/flash_attention_fwd_impl.hpp b/example/ck_tile/tutorial/03_flash_attention_fwd/flash_attention_fwd_impl.hpp index c7c7ead371..28b55d4bb4 100644 --- a/example/ck_tile/tutorial/03_flash_attention_fwd/flash_attention_fwd_impl.hpp +++ b/example/ck_tile/tutorial/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); @@ -210,7 +210,7 @@ struct FlashAttentionFwdImpl {0, 0}, make_static_tile_distribution(gemm1.MakeBBlockDistributionEncode())); #else - auto v_lds_window = make_tile_window( + auto v_lds_window = make_tile_window( v_lds, make_tuple(number{}, number{}), {0, 0}); #endif diff --git a/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp b/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp index d7e6ee4c96..9943fdbef4 100644 --- a/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/block_gemm_areg_bsmem_creg_v1_iteratek_policy.hpp +++ b/example/ck_tile/tutorial/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/tutorial/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp b/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp index 32cb98b886..3588d1dae5 100644 --- a/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/flash_attention_fwd.hpp +++ b/example/ck_tile/tutorial/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/tutorial/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp b/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp index c7c7ead371..28b55d4bb4 100644 --- a/example/ck_tile/tutorial/04_codegen_flash_attention_fwd/flash_attention_fwd_impl.hpp +++ b/example/ck_tile/tutorial/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); @@ -210,7 +210,7 @@ struct FlashAttentionFwdImpl {0, 0}, make_static_tile_distribution(gemm1.MakeBBlockDistributionEncode())); #else - auto v_lds_window = make_tile_window( + auto v_lds_window = make_tile_window( v_lds, make_tuple(number{}, number{}), {0, 0}); #endif