Add instances to consider for adding softmax support

2026-05-15 10:37:44 +00:00 · 2025-10-14 09:40:23 +00:00
parent 2072e53d1e
commit fdb89d3e2f
204 changed files with 3018 additions and 305 deletions
--- a/example/ck_tile/18_hstu_attention/generate_instances.py
+++ b/example/ck_tile/18_hstu_attention/generate_instances.py
@@ -26,16 +26,17 @@ HSTU_FORWARD_INSTANCE_TEMPLATE_INC = """
 """

 HSTU_FORWARD_INSTANCE_TEMPLATE = """
-{extern}template void run_{mode}_forward_causal_bias_dropout_dispatch<
+{extern}template void run_{mode}_forward_causal_softmax_bias_dropout_dispatch<
    {dtype},
    {has_causal},
+    {use_softmax},
    {has_bias},
    {has_dropout},
    {max_k}>(HstuAttentionFwdParams& param, hipStream_t stream);
 """

 HSTU_FORWARD_INSTANCE_FNAME = (
-    "hstu_attention_{mode}_forward_{dtype_str}_{has_or_no_causal_str}_"
+    "hstu_attention_{mode}_forward_{dtype_str}_{has_or_no_causal_str}_{use_softmax_or_not_str}_"
    "{has_or_no_bias_str}_{has_or_no_dropout_str}_{max_k_str}.cpp"
 )

@@ -48,6 +49,11 @@ BOOL_MAP_CAUSAL = {
    False: "no_causal",
 }

+BOOL_MAP_SOFTMAX = {
+    True: "softmax_true",
+    False: "softmax_false",
+}
+
 BOOL_MAP_BIAS = {
    True: "has_bias",
    False: "no_bias",
@@ -79,38 +85,41 @@ def create_forward_instances(instance_dir: Path, headdims: List) -> None:
    for mode in ["batched", "jagged"]:
        for dtype in ["fp16", "bf16"]:
            for has_causal in [True, False]:
-                for has_bias in [True, False]:
-                    for has_dropout in [True, False]:
-                        for max_k in headdims:
-                            fname = HSTU_FORWARD_INSTANCE_FNAME.format(
-                                mode=mode,
-                                dtype_str=dtype,
-                                has_or_no_causal_str=BOOL_MAP_CAUSAL[has_causal],
-                                has_or_no_bias_str=BOOL_MAP_BIAS[has_bias],
-                                has_or_no_dropout_str=BOOL_MAP_DROPOUT[has_dropout],
-                                max_k_str=INT_MAP_MAX_K[max_k],
-                            )
-                            forward_instance_inc = (
-                                HSTU_FORWARD_INSTANCE_TEMPLATE_INC.format(
+                for use_softmax in [True, False]:
+                    for has_bias in [True, False]:
+                        for has_dropout in [True, False]:
+                            for max_k in headdims:
+                                fname = HSTU_FORWARD_INSTANCE_FNAME.format(
                                    mode=mode,
-                                    dtype_file=TYPE_FNAME_MAP[dtype],
+                                    dtype_str=dtype,
+                                    has_or_no_causal_str=BOOL_MAP_CAUSAL[has_causal],
+                                    use_softmax_or_not_str=BOOL_MAP_SOFTMAX[use_softmax],
+                                    has_or_no_bias_str=BOOL_MAP_BIAS[has_bias],
+                                    has_or_no_dropout_str=BOOL_MAP_DROPOUT[has_dropout],
+                                    max_k_str=INT_MAP_MAX_K[max_k],
                                )
-                            )
-                            forward_instance = HSTU_FORWARD_INSTANCE_TEMPLATE.format(
-                                extern="",
-                                mode=mode,
-                                dtype=TYPE_CTYPE_MAP[dtype],
-                                has_causal=BOOL_MAP[has_causal],
-                                has_bias=BOOL_MAP[has_bias],
-                                has_dropout=BOOL_MAP[has_dropout],
-                                max_k=max_k,
-                                cap_mode=MODE_NAME_MAP[mode],
-                            )
-                            (instance_dir / fname).write_text(
-                                HSTU_COPYRIGHT_HEADER
-                                + forward_instance_inc
-                                + forward_instance
-                            )
+                                forward_instance_inc = (
+                                    HSTU_FORWARD_INSTANCE_TEMPLATE_INC.format(
+                                        mode=mode,
+                                        dtype_file=TYPE_FNAME_MAP[dtype],
+                                    )
+                                )
+                                forward_instance = HSTU_FORWARD_INSTANCE_TEMPLATE.format(
+                                    extern="",
+                                    mode=mode,
+                                    dtype=TYPE_CTYPE_MAP[dtype],
+                                    has_causal=BOOL_MAP[has_causal],
+                                    use_softmax=BOOL_MAP[use_softmax],
+                                    has_bias=BOOL_MAP[has_bias],
+                                    has_dropout=BOOL_MAP[has_dropout],
+                                    max_k=max_k,
+                                    cap_mode=MODE_NAME_MAP[mode],
+                                )
+                                (instance_dir / fname).write_text(
+                                     HSTU_COPYRIGHT_HEADER
+                                     + forward_instance_inc
+                                     + forward_instance
+                                )    


 def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
@@ -133,19 +142,21 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
                    for has_bias in [True, False]:
                        for has_dropout in [True, False]:
                            for has_causal in [True, False]:
-                                forward_instance = (
-                                    HSTU_FORWARD_INSTANCE_TEMPLATE.format(
-                                        extern="extern ",
-                                        mode=mode,
-                                        dtype=TYPE_CTYPE_MAP[dtype],
-                                        has_causal=BOOL_MAP[has_causal],
-                                        has_bias=BOOL_MAP[has_bias],
-                                        has_dropout=BOOL_MAP[has_dropout],
-                                        max_k=max_k,
-                                        cap_mode=MODE_NAME_MAP[mode],
-                                    )
-                                )
-                                file.write(forward_instance)
+                                for use_softmax in [True, False]:
+                                    forward_instance = (
+                                        HSTU_FORWARD_INSTANCE_TEMPLATE.format(
+                                            extern="extern ",
+                                            mode=mode,
+                                            dtype=TYPE_CTYPE_MAP[dtype],
+                                            has_causal=BOOL_MAP[has_causal],
+                                            use_softmax=BOOL_MAP[use_softmax],
+                                            has_bias=BOOL_MAP[has_bias],
+                                            has_dropout=BOOL_MAP[has_dropout],
+                                            max_k=max_k,
+                                            cap_mode=MODE_NAME_MAP[mode],
+                                        )
+                                    )   
+                                    file.write(forward_instance)

 if __name__ == "__main__":
    headdims_fwd = [64, 128, 256]
--- a/example/ck_tile/18_hstu_attention/hstu_attention_batched_forward_bf16.cpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_batched_forward_bf16.cpp
@@ -17,11 +17,12 @@ void hstu_attention_batched_forward_bf16(HstuAttentionFwdParams& param, hipStrea
    const bool use_causal  = param.use_causal;
    BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
        HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
-            run_batched_forward_causal_bias_dropout_dispatch<ck_tile::bf16_t,
-                                                             kUseCausal,
-                                                             kHasBias,
-                                                             kHasDropout,
-                                                             MaxK>(param, stream);
+            run_batched_forward_causal_softmax_bias_dropout_dispatch<ck_tile::bf16_t,
+                                                                     kUseCausal,
+                                                                     false, // using softmax
+                                                                     kHasBias,
+                                                                     kHasDropout,
+                                                                     MaxK>(param, stream);
        });
    });
 };
--- a/example/ck_tile/18_hstu_attention/hstu_attention_batched_forward_dispatch.hpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_batched_forward_dispatch.hpp
@@ -21,10 +21,11 @@

 template <typename InOutDataType,
          bool kUseCausal,
+          bool kUseSoftmax,
          bool kHasBias,
          bool kHasDropout,
          ck_tile::index_t MaxK>
-struct batched_forward_causal_bias_dropout_dispatch
+struct batched_forward_causal_softmax_bias_dropout_dispatch
 {
    using HstuAttentionTileSetting = typename HstuAttentionFwdTileSetting<MaxK>::Type;

@@ -38,6 +39,7 @@ struct batched_forward_causal_bias_dropout_dispatch
        kHasBias,
        kHasDropout,
        kUseCausal,
+        kUseSoftmax,
        HstuAttentionTileSetting,
        HstuTraits>;

@@ -137,15 +139,17 @@ struct batched_forward_causal_bias_dropout_dispatch

 template <typename InOutDataType,
          bool kUseCausal,
+          bool kUseSoftmax,
          bool kHasBias,
          bool kHasDropout,
          ck_tile::index_t MaxK>
-void run_batched_forward_causal_bias_dropout_dispatch(HstuAttentionFwdParams& param,
-                                                      hipStream_t stream)
+void run_batched_forward_causal_softmax_bias_dropout_dispatch(HstuAttentionFwdParams& param,
+                                                              hipStream_t stream)
 {
-    batched_forward_causal_bias_dropout_dispatch<InOutDataType,
-                                                 kUseCausal,
-                                                 kHasBias,
-                                                 kHasDropout,
-                                                 MaxK>::Run(param, stream);
+    batched_forward_causal_softmax_bias_dropout_dispatch<InOutDataType,
+                                                         kUseCausal,
+                                                         kUseSoftmax,
+                                                         kHasBias,
+                                                         kHasDropout,
+                                                         MaxK>::Run(param, stream);
 };
--- a/example/ck_tile/18_hstu_attention/hstu_attention_batched_forward_fp16.cpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_batched_forward_fp16.cpp
@@ -17,11 +17,12 @@ void hstu_attention_batched_forward_fp16(HstuAttentionFwdParams& param, hipStrea
    const bool use_causal  = param.use_causal;
    BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
        HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
-            run_batched_forward_causal_bias_dropout_dispatch<ck_tile::fp16_t,
-                                                             kUseCausal,
-                                                             kHasBias,
-                                                             kHasDropout,
-                                                             MaxK>(param, stream);
+            run_batched_forward_causal_softmax_bias_dropout_dispatch<ck_tile::fp16_t,
+                                                                     kUseCausal,
+                                                                     false, // using softmax
+                                                                     kHasBias,
+                                                                     kHasDropout,
+                                                                     MaxK>(param, stream);
        });
    });
 };
--- a/example/ck_tile/18_hstu_attention/hstu_attention_jagged_forward_bf16.cpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_jagged_forward_bf16.cpp
@@ -17,11 +17,12 @@ void hstu_attention_jagged_forward_bf16(HstuAttentionFwdParams& param, hipStream
    const bool use_causal  = param.use_causal;
    BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
        HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
-            run_jagged_forward_causal_bias_dropout_dispatch<ck_tile::bf16_t,
-                                                            kUseCausal,
-                                                            kHasBias,
-                                                            kHasDropout,
-                                                            MaxK>(param, stream);
+            run_jagged_forward_causal_softmax_bias_dropout_dispatch<ck_tile::bf16_t,
+                                                                    kUseCausal,
+                                                                    false, // using softmax
+                                                                    kHasBias,
+                                                                    kHasDropout,
+                                                                    MaxK>(param, stream);
        });
    });
 };
--- a/example/ck_tile/18_hstu_attention/hstu_attention_jagged_forward_dispatch.hpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_jagged_forward_dispatch.hpp
@@ -21,10 +21,11 @@

 template <typename InOutDataType,
          bool kUseCausal,
+          bool kUseSoftmax,
          bool kHasBias,
          bool kHasDropout,
          ck_tile::index_t MaxK>
-struct jagged_forward_causal_bias_dropout_dispatch
+struct jagged_forward_causal_softmax_bias_dropout_dispatch
 {
    using HstuAttentionTileSetting = typename HstuAttentionFwdTileSetting<MaxK>::Type;

@@ -38,6 +39,7 @@ struct jagged_forward_causal_bias_dropout_dispatch
        kHasBias,
        kHasDropout,
        kUseCausal,
+        kUseSoftmax,
        HstuAttentionTileSetting,
        HstuTraits>;

@@ -128,15 +130,17 @@ struct jagged_forward_causal_bias_dropout_dispatch

 template <typename InOutDataType,
          bool kUseCausal,
+          bool kUseSoftmax,
          bool kHasBias,
          bool kHasDropout,
          ck_tile::index_t MaxK>
-void run_jagged_forward_causal_bias_dropout_dispatch(HstuAttentionFwdParams& param,
-                                                     hipStream_t stream)
+void run_jagged_forward_causal_softmax_bias_dropout_dispatch(HstuAttentionFwdParams& param,
+                                                             hipStream_t stream)
 {
-    jagged_forward_causal_bias_dropout_dispatch<InOutDataType,
-                                                kUseCausal,
-                                                kHasBias,
-                                                kHasDropout,
-                                                MaxK>::Run(param, stream);
+    jagged_forward_causal_softmax_bias_dropout_dispatch<InOutDataType,
+                                                        kUseCausal,
+                                                        kUseSoftmax,
+                                                        kHasBias,
+                                                        kHasDropout,
+                                                        MaxK>::Run(param, stream);
 };
--- a/example/ck_tile/18_hstu_attention/hstu_attention_jagged_forward_fp16.cpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_jagged_forward_fp16.cpp
@@ -17,11 +17,12 @@ void hstu_attention_jagged_forward_fp16(HstuAttentionFwdParams& param, hipStream
    const bool use_causal  = param.use_causal;
    BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
        HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
-            run_jagged_forward_causal_bias_dropout_dispatch<ck_tile::fp16_t,
-                                                            kUseCausal,
-                                                            kHasBias,
-                                                            kHasDropout,
-                                                            MaxK>(param, stream);
+            run_jagged_forward_causal_softmax_bias_dropout_dispatch<ck_tile::fp16_t,
+                                                                    kUseCausal,
+                                                                    false, // using softmax
+                                                                    kHasBias,
+                                                                    kHasDropout,
+                                                                    MaxK>(param, stream);
        });
    });
 };
--- a/example/ck_tile/18_hstu_attention/hstu_attention_pipeline_problem.hpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_pipeline_problem.hpp
@@ -21,6 +21,7 @@ template <typename InOutDataType_,
          bool kHasBias_,
          bool kHasDropout_,
          bool kHasCausal_,
+          bool kUseSoftmax_,
          typename AttentionTileSetting_,
          typename Traits_>
 struct HstuAttentionFwdPipelineProblem
@@ -42,6 +43,7 @@ struct HstuAttentionFwdPipelineProblem
    static constexpr bool kHasBias    = kHasBias_;
    static constexpr bool kHasDropout = kHasDropout_;
    static constexpr bool kHasCausal  = kHasCausal_;
+    static constexpr bool kUseSoftmax = kUseSoftmax_;

    using HstuAttentionTileSetting = remove_cvref_t<AttentionTileSetting_>;

--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    false,
    true,
+    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
-    true,
    false,
+    true,
+    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    false,
    true,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    true,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    true,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    true,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    true,
    true,
+    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    true,
    true,
+    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    true,
    true,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
+    true,
+    true,
    false,
-    true,
-    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
-    false,
    true,
+    true,
+    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    true,
+    true,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    true,
    false,
+    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
+    true,
    false,
    true,
-    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
+    true,
    false,
    true,
-    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_no_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
+    true,
    false,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_no_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
+    true,
    false,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_no_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
+    true,
    false,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp
@@ -9,169 +9,385 @@
 #include <ck_tile/core/numeric/half.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
    true,
    true,
    true,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    true,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    false,
    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    false,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    true,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    true,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    false,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    false,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    true,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    true,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    true,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    false,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
    false,
    false,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    true,
    true,
    false,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::bf16_t,
+    true,
+    false,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    false,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/half.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::bf16_t,
+    false,
+    true,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    false,
    true,
+    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
-    true,
    false,
+    true,
+    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    false,
    true,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    true,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    true,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    false,
    false,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
@@ -9,8 +9,9 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    true,
    true,
+    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    true,
    true,
+    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    true,
    true,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
+    true,
+    true,
    false,
-    true,
-    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
-    false,
    true,
+    true,
+    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    true,
+    true,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    true,
    false,
+    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
+    true,
    false,
    true,
-    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
+    true,
    false,
    true,
-    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_no_dropout_maxk_128.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
+    true,
    false,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_no_dropout_maxk_256.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
+    true,
    false,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_no_dropout_maxk_64.cpp
@@ -9,9 +9,10 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-template void run_batched_forward_causal_bias_dropout_dispatch<
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
-    false,
-    false,
    true,
+    true,
+    false,
+    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp
@@ -9,169 +9,385 @@
 #include <ck_tile/core/numeric/bfloat16.hpp>
 #include "hstu_attention_batched_forward_dispatch.hpp"

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
    true,
    true,
    true,
+    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    true,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    false,
    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    true,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    false,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    false,
    64>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    true,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    true,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    false,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    true,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    false,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    false,
    128>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    true,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    true,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    true,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    false,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
    false,
    false,
    true,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    false,
+    true,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    true,
    true,
    false,
    false,
    256>(HstuAttentionFwdParams& param, hipStream_t stream);

-extern template void run_batched_forward_causal_bias_dropout_dispatch<
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
    ck_tile::fp16_t,
+    true,
+    false,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
+
+extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
    false,
    false,
    false,
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    false,
+    false,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    true,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    true,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    true,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_no_dropout_maxk_128.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    false,
+    128>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_no_dropout_maxk_256.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    false,
+    256>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
+++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_no_dropout_maxk_64.cpp
@@ -0,0 +1,18 @@
+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+// The file is automatically generated, don't modify!
+// See the generator script
+// `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py`
+
+#include <ck_tile/core/numeric/bfloat16.hpp>
+#include "hstu_attention_batched_forward_dispatch.hpp"
+
+template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
+    ck_tile::fp16_t,
+    false,
+    true,
+    true,
+    false,
+    64>(HstuAttentionFwdParams& param, hipStream_t stream);
--- a/Show More
+++ b/Show More