From 179f0e857e92ac33ef71fcfe53532ffe0178ee4c Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Sun, 14 Dec 2025 16:21:54 +0000 Subject: [PATCH] Rename WarpTile in fwd setting --- .../hstu_attention_fwd_setting.hpp | 78 +++++++++---------- 1 file changed, 39 insertions(+), 39 deletions(-) diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_setting.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_setting.hpp index 50889863d5..2f6844c483 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_setting.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_setting.hpp @@ -8,9 +8,9 @@ #include "hstu_attention_fwd_type_config.hpp" #include "hstu_attention_tile_setting_define.hpp" -using HstuAttentionFwdWarpTile1 = ck_tile::sequence<16, 16, 16>; -using HstuAttentionFwdWarpTile2 = ck_tile::sequence<16, 16, 32>; -using HstuAttentionFwdWarpTile3 = ck_tile::sequence<32, 32, 16>; +using WarpTile_16x16x16 = ck_tile::sequence<16, 16, 16>; +using WarpTile_16x16x32 = ck_tile::sequence<16, 16, 32>; +using WarpTile_32x32x16 = ck_tile::sequence<32, 32, 16>; #if !defined(BUILD_HSTU_FOR_GFX95_ONLY) template @@ -96,9 +96,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<64> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<64>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<64>::gemm0_warps, - HstuAttentionFwdWarpTile1, + WarpTile_16x16x16, typename HstuAttentionNoSoftmaxFwdBlockTile<64>::gemm1_warps, - HstuAttentionFwdWarpTile1>; + WarpTile_16x16x16>; }; template <> @@ -107,9 +107,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<96> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<96>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<96>::gemm0_warps, - HstuAttentionFwdWarpTile1, + WarpTile_16x16x16, typename HstuAttentionNoSoftmaxFwdBlockTile<96>::gemm1_warps, - HstuAttentionFwdWarpTile1>; + WarpTile_16x16x16>; }; template <> @@ -118,9 +118,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<128> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<128>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<128>::gemm0_warps, - HstuAttentionFwdWarpTile1, + WarpTile_16x16x16, typename HstuAttentionNoSoftmaxFwdBlockTile<128>::gemm1_warps, - HstuAttentionFwdWarpTile1>; + WarpTile_16x16x16>; }; template <> @@ -129,9 +129,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<256> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<256>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<256>::gemm0_warps, - HstuAttentionFwdWarpTile1, + WarpTile_16x16x16, typename HstuAttentionNoSoftmaxFwdBlockTile<256>::gemm1_warps, - HstuAttentionFwdWarpTile1>; + WarpTile_16x16x16>; }; template @@ -143,9 +143,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<64> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<64>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<64>::gemm0_warps, - HstuAttentionFwdWarpTile3, + WarpTile_32x32x16, typename HstuAttentionWithSoftmaxFwdBlockTile<64>::gemm1_warps, - HstuAttentionFwdWarpTile3>; + WarpTile_32x32x16>; }; template <> @@ -154,9 +154,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<96> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<96>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<96>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_32x32x16, typename HstuAttentionWithSoftmaxFwdBlockTile<96>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_32x32x16>; }; template <> @@ -165,9 +165,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<128> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<128>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<128>::gemm0_warps, - HstuAttentionFwdWarpTile1, + WarpTile_16x16x16, typename HstuAttentionWithSoftmaxFwdBlockTile<128>::gemm1_warps, - HstuAttentionFwdWarpTile1>; + WarpTile_16x16x16>; }; template <> @@ -176,9 +176,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<256> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<256>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<256>::gemm0_warps, - HstuAttentionFwdWarpTile1, + WarpTile_16x16x16, typename HstuAttentionWithSoftmaxFwdBlockTile<256>::gemm1_warps, - HstuAttentionFwdWarpTile1>; + WarpTile_16x16x16>; }; #endif @@ -282,9 +282,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<32> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<32>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<32>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionNoSoftmaxFwdBlockTile<32>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template <> @@ -293,9 +293,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<64> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<64>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<64>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionNoSoftmaxFwdBlockTile<64>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template <> @@ -304,9 +304,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<96> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<96>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<96>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionNoSoftmaxFwdBlockTile<96>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template <> @@ -315,9 +315,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<128> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<128>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<128>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionNoSoftmaxFwdBlockTile<128>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template <> @@ -326,9 +326,9 @@ struct HstuAttentionNoSoftmaxFwdTileSetting<256> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionNoSoftmaxFwdBlockTile<256>::type, typename HstuAttentionNoSoftmaxFwdBlockTile<256>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionNoSoftmaxFwdBlockTile<256>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template @@ -340,9 +340,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<32> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<32>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<32>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionWithSoftmaxFwdBlockTile<32>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template <> @@ -351,9 +351,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<64> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<64>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<64>::gemm0_warps, - HstuAttentionFwdWarpTile3, + WarpTile_32x32x16, typename HstuAttentionWithSoftmaxFwdBlockTile<64>::gemm1_warps, - HstuAttentionFwdWarpTile3>; + WarpTile_32x32x16>; }; template <> @@ -362,9 +362,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<96> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<96>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<96>::gemm0_warps, - HstuAttentionFwdWarpTile3, + WarpTile_32x32x16, typename HstuAttentionWithSoftmaxFwdBlockTile<96>::gemm1_warps, - HstuAttentionFwdWarpTile3>; + WarpTile_32x32x16>; }; template <> @@ -373,9 +373,9 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<128> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<128>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<128>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionWithSoftmaxFwdBlockTile<128>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; template <> @@ -384,8 +384,8 @@ struct HstuAttentionWithSoftmaxFwdTileSetting<256> using Type = ck_tile::HstuAttentionFwdTileSettingClass< typename HstuAttentionWithSoftmaxFwdBlockTile<256>::type, typename HstuAttentionWithSoftmaxFwdBlockTile<256>::gemm0_warps, - HstuAttentionFwdWarpTile2, + WarpTile_16x16x32, typename HstuAttentionWithSoftmaxFwdBlockTile<256>::gemm1_warps, - HstuAttentionFwdWarpTile2>; + WarpTile_16x16x32>; }; #endif