ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-17 11:30:02 +00:00
Code Issues Packages Projects Releases Wiki Activity

[Performance] Use N0Sub=16 for trload with softmax pipeline to reduce vgpr spilling

Browse Source
This commit is contained in:
Qianfeng Zhang
2026-02-02 15:59:38 +00:00
parent c360e0cbc4
commit 0a8c5f523a
1 changed files with 2 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
example/ck_tile/18_hstu_attention/hstu_attention_fwd_setting.hpp
View File

@@ -260,7 +260,7 @@ struct HstuAttentionWithSoftmaxFwdBlockTile<96>
template <>
struct HstuAttentionWithSoftmaxFwdBlockTile<128>
{
using type = ck_tile::sequence<128, 64, 32, 128, 32, 128>;
using type = ck_tile::sequence<128, 64, 16, 128, 32, 128>;
using gemm0_warps = ck_tile::sequence<4, 1, 1>;
using gemm1_warps = ck_tile::sequence<4, 1, 1>;
};
@@ -268,7 +268,7 @@ struct HstuAttentionWithSoftmaxFwdBlockTile<128>
template <>
struct HstuAttentionWithSoftmaxFwdBlockTile<256>
{
using type = ck_tile::sequence<128, 64, 32, 256, 32, 256>;
using type = ck_tile::sequence<128, 64, 16, 256, 32, 256>;
using gemm0_warps = ck_tile::sequence<4, 1, 1>;
using gemm1_warps = ck_tile::sequence<4, 1, 1>;
};