ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-13 18:51:13 +00:00
Code Issues Packages Projects Releases Wiki Activity

Kernel use types declared in the problem rather than the pipeline

Browse Source
This commit is contained in:
Qianfeng Zhang
2026-06-04 08:13:24 +00:00
parent ee5bd0ebba
commit fd9af72c9f
1 changed files with 5 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp
View File

@@ -36,9 +36,11 @@ struct HstuAttentionFwdKernel
static constexpr ck_tile::index_t kBlockPerCu = HstuAttentionPipeline::kBlockPerCu;
static_assert(kBlockPerCu > 0);
using QKVDataType = ck_tile::remove_cvref_t<typename HstuAttentionPipeline::QKVDataType>;
using BiasDataType = ck_tile::remove_cvref_t<typename HstuAttentionPipeline::BiasDataType>;
using ODataType = ck_tile::remove_cvref_t<typename HstuAttentionPipeline::ODataType>;
using QKVDataType =
ck_tile::remove_cvref_t<typename HstuAttentionPipeline::Problem::QKVDataType>;
using BiasDataType =
ck_tile::remove_cvref_t<typename HstuAttentionPipeline::Problem::BiasDataType>;
using ODataType = ck_tile::remove_cvref_t<typename HstuAttentionPipeline::Problem::ODataType>;
static constexpr bool kIsCrossAttention = HstuAttentionPipeline::Problem::kIsCrossAttention;
static constexpr bool kUseGroup = HstuAttentionPipeline::Problem::kUseGroup;