From fd9af72c9f5cb6feb390994c5f286ffda3dd8010 Mon Sep 17 00:00:00 2001
From: Qianfeng Zhang <Qianfeng.Zhang@amd.com>
Date: Thu, 4 Jun 2026 08:13:24 +0000
Subject: [PATCH] Kernel use types declared in the problem rather than the
 pipeline

---
 .../18_hstu_attention/hstu_attention_fwd_kernel.hpp       | 8 +++++---
 1 file changed, 5 insertions(+), 3 deletions(-)

diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp
index bb9977d189..2b9f7a3d2f 100644
--- a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp
@@ -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;